- a patch to be able to load the CUDA library at runtime and only when available.
authorsgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Mon, 2 Nov 2009 07:19:00 +0000 (07:19 +0000)
committersgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Mon, 2 Nov 2009 07:19:00 +0000 (07:19 +0000)
- cleanup

53 files changed:
HLT/TPCLib/tracking-ca/AliHLTTPCCAClusterData.h
HLT/TPCLib/tracking-ca/AliHLTTPCCADataCompressor.h
HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h
HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cu.patch [new file with mode: 0644]
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cu
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cxx [new file with mode: 0644]
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.h [new file with mode: 0644]
HLT/TPCLib/tracking-ca/AliHLTTPCCAGlobalMergerComponent.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMCPoint.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMergedTrack.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMerger.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMergerOutput.h
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.h
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAOutTrack.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAPerformance.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h
HLT/TPCLib/tracking-ca/AliHLTTPCCARow.h
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceData.h
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceOutput.h
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceTrack.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAStandaloneFramework.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsSorter.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrack.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackLinearisation.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam2.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerFramework.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerFramework.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATracklet.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.h
HLT/TPCLib/tracking-ca/AliTPCtrackerCA.h
HLT/TPCLib/tracking-ca/MemoryAssignmentHelpers.h
HLT/libAliHLTTPC.conf [deleted file]

index df9c1f6..20c6353 100644 (file)
@@ -152,11 +152,7 @@ class AliHLTTPCCAClusterData
      */
     void Merge( int index1, int index2 );
 
-#ifdef REPRODUCIBLE_CLUSTER_SORTING
-       static bool CompareClusters( const Data &a, const Data &b ) { return ( a.fRow >= b.fRow ? (a.fId < b.fId) : (a.fRow < b.fRow) ); }
-#else
-       static bool CompareClusters( const Data &a, const Data &b ) { return ( (a.fRow < b.fRow) ); }
-#endif
+       static bool CompareClusters( const Data &a, const Data &b ) { return ( a.fRow == b.fRow ? (a.fY < b.fY) : (a.fRow < b.fRow) ); }
 
     int fSliceIndex;  // the slice index this data belongs to
     int fFirstRow; // see FirstRow()
index 9b6800e..cd48cb9 100644 (file)
@@ -79,4 +79,4 @@ GPUhd() inline float AliHLTTPCCADataCompressor::UShort2Z( unsigned short iYZ )
   return ( iYZ % 256 )*kMult - 3.f;
 }
 
-#endif
+#endif //ALIHLTTPCCADATACOMPRESSOR_H
index 5e37b56..0ea7a57 100644 (file)
 
 #ifdef __CUDACC__
 #define HLTCA_GPUCODE
-#endif
+#endif //__CUDACC__
 
 #ifdef WIN32
 #ifndef R__WIN32
 #define R__WIN32
-#endif
-#endif
+#endif //!R__WIN32
+#endif //WIN32
 
-#if defined(R__WIN32)
+#ifdef R__WIN32
 #ifdef INTEL_RUNTIME
 #pragma warning(disable : 1786)
 #pragma warning(disable : 1478)
 #pragma warning(disable : 161)
-#endif
+#endif //INTEL_RUNTIME
 
 #ifdef VSNET_RUNTIME
 #pragma warning(disable : 4616)
 #pragma warning(disable : 4996)
 #pragma warning(disable : 1684)
-#endif
-#endif
+#endif //VSNET_RUNTIME
+#endif //R__WIN32
 
-#if defined(HLTCA_STANDALONE)
+#ifdef HLTCA_STANDALONE
 
 // class TObject{};
 
 #ifdef ClassDef
 #undef ClassDef
-#endif
+#endif //ClassDef
 
 #ifdef ClassTmp
 #undef ClassTmp
-#endif
+#endif //ClassTmp
 
 #define ClassDef(name,id)
 #define ClassImp(name)
@@ -64,7 +64,7 @@ typedef unsigned long  ULong_t;     //Unsigned long integer 8 bytes (unsigned lo
 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
+#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)
@@ -79,7 +79,7 @@ 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)
@@ -103,11 +103,12 @@ namespace AliHLTTPCCADefinitions
   extern const AliHLTComponentDataType fgkCompressedInputDataType;
 }
 
-#endif
+#endif //HLTCA_STANDALONE
 
 //#define EXTERN_ROW_HITS
 #define TRACKLET_SELECTOR_MIN_HITS 10
 #define REPRODUCIBLE_CLUSTER_SORTING
+//#define FAST_NEIGHBOURS_FINDER
 
 #ifdef HLTCA_GPUCODE
 #define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 6
@@ -119,7 +120,7 @@ namespace AliHLTTPCCADefinitions
 #define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 7000
 #define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 10000
 #define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 15000
-#endif
+#endif //HLTCA_GPUCODE
 
 #ifdef HLTCA_GPUCODE
 
@@ -159,9 +160,9 @@ inline bool finite(float x)
 {
        return(x <= FLT_MAX);
 }
-#endif
+#endif //R__WIN32
 
-#endif
+#endif //HLTCA_GPUCODE
 
 /*
  * Helper for compile-time verification of correct API usage
@@ -182,7 +183,7 @@ namespace
   (void) Error_##msg
 #else
 #define STATIC_ASSERT(a, b)
-#endif
+#endif //!HLTCA_GPUCODE
 
 namespace
 {
@@ -212,4 +213,4 @@ namespace
 #define UNROLL16(var, code) UNROLL8(var, code) UNROLL8(var, code)
 #define UNROLL32(var, code) UNROLL16(var, code) UNROLL16(var, code)
 
-#endif
+#endif //ALIHLTTPCCADEF_H
index 5011e1c..37d2ce2 100644 (file)
@@ -112,4 +112,4 @@ class AliHLTTPCCADisplay
 
 };
 
-#endif
+#endif //ALIHLTTPCCADISPLAY_H
diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cu.patch b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cu.patch
new file mode 100644 (file)
index 0000000..de94951
--- /dev/null
@@ -0,0 +1,11 @@
+--- AliHLTTPCCAGPUTracker.cucpp        2009-05-28 12:14:09.000000000 +0200
++++ release/x86_64-pc-linux-gnu/code/AliHLTTPCCAGPUTracker.cucpp       2009-05-28 12:10:25.000000000 +0200
+@@ -23186,7 +23186,7 @@
+ static T2 *Alloc(int s) { auto T2 *p = (reinterpret_cast< T2 *>(_mm_malloc(s * sizeof(CacheLineSizeHelper< T> ), 128))); return new (p) T2 [s]; }
+ static void Free(T2 *const p, int size) {
+ for (int i = 0; i < size; ++i) {
+-((p[i]).~CacheLineSizeHelper());
++((p[i]).~T2());
+ }
+ _mm_free(p);
+ }
index cb2a3f8..1c94411 100644 (file)
 //If not building GPU Code then build dummy functions to link against
 #include "AliHLTTPCCAGPUTracker.h"
 
-#ifndef BUILD_GPU
-int AliHLTTPCCAGPUTracker::InitGPU(int /*sliceCount*/, int /*forceDeviceID*/)
-{
-    //Dummy init function if CUDA is not available
-    HLTInfo("CUDA Compiler was not available during build process, omitting CUDA initialization");
-    return(1);
-}
-void AliHLTTPCCAGPUTracker::StandalonePerfTime(int /*iSlice*/, int /*i*/) {}
-template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int alignment) {return(NULL);}
-//bool AliHLTTPCCAGPUTracker::CudaFailedMsg(cudaError_t error) {return(true);}
-int AliHLTTPCCAGPUTracker::CUDASync(char* /*text*/) {return(1);}
+AliHLTTPCCAGPUTracker::AliHLTTPCCAGPUTracker() {}
+AliHLTTPCCAGPUTracker::~AliHLTTPCCAGPUTracker() {}
+int AliHLTTPCCAGPUTracker::InitGPU(int /*sliceCount*/, int /*forceDeviceID*/) { return(1); }
 void AliHLTTPCCAGPUTracker::SetDebugLevel(int /*dwLevel*/, std::ostream* /*NewOutFile*/) {}
 int AliHLTTPCCAGPUTracker::SetGPUTrackerOption(char* /*OptionName*/, int /*OptionValue*/) {return(1);}
 int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCASliceOutput** /*pTracker*/, AliHLTTPCCAClusterData* /*pClusterData*/, int /*fFirstSlice*/, int /*fSliceCount*/) {return(1);}
 int AliHLTTPCCAGPUTracker::ExitGPU() {return(0);}
 int AliHLTTPCCAGPUTracker::InitializeSliceParam(int /*iSlice*/, AliHLTTPCCAParam& /*param*/) { return 1; }
-void AliHLTTPCCAGPUTracker::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* /*val*/) {};
-int AliHLTTPCCAGPUTracker::GetThread(){ return 0; }
-void AliHLTTPCCAGPUTracker::ReleaseGlobalLock(void* /*sem*/) {};
-int AliHLTTPCCAGPUTracker::CheckMemorySizes(int /*sliceCount*/){ return(1); }
-#endif
+void AliHLTTPCCAGPUTracker::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* /*val*/) {}
+unsigned long long int* AliHLTTPCCAGPUTracker::PerfTimer(int /*iSlice*/, unsigned int /*i*/)
+{
+       //Just return some dummy adress the caller can access
+       static unsigned long long int tmp;
+       return(&tmp);
+}
+const AliHLTTPCCASliceOutput::outputControlStruct* AliHLTTPCCAGPUTracker::OutputControl() const { return NULL; }
+int AliHLTTPCCAGPUTracker::GetSliceCount() const { return(0); }
\ No newline at end of file
index 6d5f4ed..fdbbaf6 100644 (file)
@@ -9,93 +9,38 @@
 #define ALIHLTTPCCAGPUTRACKER_H
 
 #include "AliHLTTPCCADef.h"
-#include "AliHLTTPCCATracker.h"
-#include "AliHLTLogging.h"
 #include "AliHLTTPCCASliceOutput.h"
+#include <iostream>
 
-class AliHLTTPCCARow;
+class AliHLTTPCCAClusterData;
+class AliHLTTPCCASliceOutput;
+class AliHLTTPCCAParam;
 
-class AliHLTTPCCAGPUTracker : AliHLTLogging
+class AliHLTTPCCAGPUTracker
 {
 public:
-       AliHLTTPCCAGPUTracker() :
-         fGpuTracker(NULL),
-         fGPUMemory(NULL),
-         fHostLockedMemory(NULL),
-         fDebugLevel(0),
-         fOutFile(NULL),
-         fGPUMemSize(0),
-         fpCudaStreams(NULL),
-         fSliceCount(0),
-         fOutputControl(NULL),
-         fThreadId(0),
-         fCudaInitialized(0)
-         {};
-         ~AliHLTTPCCAGPUTracker() {};
+       AliHLTTPCCAGPUTracker();
+       virtual ~AliHLTTPCCAGPUTracker();
 
-       int InitGPU(int sliceCount = 12, int forceDeviceID = -1);
-       int Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTTPCCAClusterData* pClusterData, int fFirstSlice, int fSliceCount = -1);
-       int ExitGPU();
+       virtual int InitGPU(int sliceCount = 12, int forceDeviceID = -1);
+       virtual int Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTTPCCAClusterData* pClusterData, int fFirstSlice, int fSliceCount = -1);
+       virtual int ExitGPU();
 
-       void SetDebugLevel(const int dwLevel, std::ostream* const NewOutFile = NULL);
-       int SetGPUTrackerOption(char* OptionName, int OptionValue);
+       virtual void SetDebugLevel(const int dwLevel, std::ostream* const NewOutFile = NULL);
+       virtual int SetGPUTrackerOption(char* OptionName, int OptionValue);
 
-       unsigned long long int* PerfTimer(int iSlice, unsigned int i) {return(fSlaveTrackers ? fSlaveTrackers[iSlice].PerfTimer(i) : NULL); }
+       virtual unsigned long long int* PerfTimer(int iSlice, unsigned int i);
 
-       int InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param);
+       virtual int InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param);
+       virtual void SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val);
 
-       const AliHLTTPCCASliceOutput::outputControlStruct* OutputControl() const { return fOutputControl; }
-       void SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val);
-       
-       int GetSliceCount() const { return(fSliceCount); }
+       virtual const AliHLTTPCCASliceOutput::outputControlStruct* OutputControl() const;
+       virtual int GetSliceCount() const;
 
 private:
-       static void* RowMemory(void* const BaseMemory, int iSlice) { return( ((char*) BaseMemory) + iSlice * sizeof(AliHLTTPCCARow) * (HLTCA_ROW_COUNT + 1) ); }
-       static void* CommonMemory(void* const BaseMemory, int iSlice) { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + iSlice * AliHLTTPCCATracker::CommonMemorySize() ); }
-       static void* SliceDataMemory(void* const BaseMemory, int iSlice) { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + iSlice * HLTCA_GPU_SLICE_DATA_MEMORY ); }
-       void* GlobalMemory(void* const BaseMemory, int iSlice) const { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + fSliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY) + iSlice * HLTCA_GPU_GLOBAL_MEMORY ); }
-       void* TracksMemory(void* const BaseMemory, int iSlice) const { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + fSliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY) + iSlice * HLTCA_GPU_TRACKS_MEMORY ); }
-       void* TrackerMemory(void* const BaseMemory, int iSlice) const { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + fSliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY + HLTCA_GPU_TRACKS_MEMORY) + iSlice * sizeof(AliHLTTPCCATracker) ); }
-
-       void DumpRowBlocks(AliHLTTPCCATracker* tracker, int iSlice, bool check = true);
-       int GetThread();
-       void ReleaseGlobalLock(void* sem);
-       int CheckMemorySizes(int sliceCount);
-
-       AliHLTTPCCATracker *fGpuTracker;
-       void* fGPUMemory;
-       void* fHostLockedMemory;
-
-       int CUDASync(char* state = "UNKNOWN");
-       template <class T> T* alignPointer(T* ptr, int alignment);
-
-       void StandalonePerfTime(int iSlice, int i);
-
-       int fDebugLevel;                        //Debug Level for GPU Tracker
-       std::ostream* fOutFile;         //Debug Output Stream Pointer
-       unsigned long long int fGPUMemSize;     //Memory Size to allocate on GPU
-
-       void* fpCudaStreams;
-
-       int fSliceCount;
-
-       static const int fgkNSlices = 36;
-       AliHLTTPCCATracker fSlaveTrackers[fgkNSlices];
-#ifdef HLTCA_GPUCODE
-       bool CudaFailedMsg(cudaError_t error);
-#endif
-
-       AliHLTTPCCASliceOutput::outputControlStruct* fOutputControl;
-       
-       static bool fgGPUUsed;
-       int fThreadId;
-       int fCudaInitialized;
-
        // disable copy
        AliHLTTPCCAGPUTracker( const AliHLTTPCCAGPUTracker& );
        AliHLTTPCCAGPUTracker &operator=( const AliHLTTPCCAGPUTracker& );
-
-       ClassDef( AliHLTTPCCAGPUTracker, 0 )
 };
 
-#endif
+#endif //ALIHLTTPCCAGPUTRACKER_H
index 9d272d5..0713bec 100644 (file)
 //                                                                          *
 //***************************************************************************
 
-#include "AliHLTTPCCAGPUTracker.h"
-
-#ifdef BUILD_GPU
+#include "AliHLTTPCCAGPUTrackerNVCC.h"
 
+#ifdef HLTCA_GPUCODE
 #include <cuda.h>
+#include <sm_11_atomic_functions.h>
+#include <sm_12_atomic_functions.h>
+#endif
+
 #ifdef R__WIN32
 #else
 #include <sys/syscall.h>
@@ -32,8 +35,6 @@
 #include "AliHLTTPCCADef.h"
 #include "AliHLTTPCCAGPUConfig.h"
 
-#include <sm_11_atomic_functions.h>
-#include <sm_12_atomic_functions.h>
 
 #include <iostream>
 
@@ -77,13 +78,29 @@ texture<signed short, 1, cudaReadModeElementType> gAliTexRefs;
 #include "AliHLTSystem.h"
 #endif
 
-ClassImp( AliHLTTPCCAGPUTracker )
+ClassImp( AliHLTTPCCAGPUTrackerNVCC )
 
-bool AliHLTTPCCAGPUTracker::fgGPUUsed = false;
+bool AliHLTTPCCAGPUTrackerNVCC::fgGPUUsed = false;
 
 #define SemLockName "AliceHLTTPCCAGPUTrackerInitLockSem"
 
-void AliHLTTPCCAGPUTracker::ReleaseGlobalLock(void* sem)
+AliHLTTPCCAGPUTrackerNVCC::AliHLTTPCCAGPUTrackerNVCC() :
+       fGpuTracker(NULL),
+       fGPUMemory(NULL),
+       fHostLockedMemory(NULL),
+       fDebugLevel(0),
+       fOutFile(NULL),
+       fGPUMemSize(0),
+       fpCudaStreams(NULL),
+       fSliceCount(0),
+       fOutputControl(NULL),
+       fThreadId(0),
+       fCudaInitialized(0)
+       {};
+
+AliHLTTPCCAGPUTrackerNVCC::~AliHLTTPCCAGPUTrackerNVCC() {};
+
+void AliHLTTPCCAGPUTrackerNVCC::ReleaseGlobalLock(void* sem)
 {
        //Release the global named semaphore that locks GPU Initialization
 #ifdef R__WIN32
@@ -98,7 +115,7 @@ void AliHLTTPCCAGPUTracker::ReleaseGlobalLock(void* sem)
 #endif
 }
 
-int AliHLTTPCCAGPUTracker::CheckMemorySizes(int sliceCount)
+int AliHLTTPCCAGPUTrackerNVCC::CheckMemorySizes(int sliceCount)
 {
        //Check constants for correct memory sizes
   if (sizeof(AliHLTTPCCATracker) * sliceCount > HLTCA_GPU_TRACKER_OBJECT_MEMORY)
@@ -121,7 +138,7 @@ int AliHLTTPCCAGPUTracker::CheckMemorySizes(int sliceCount)
   return(0);
 }
 
-int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
+int AliHLTTPCCAGPUTrackerNVCC::InitGPU(int sliceCount, int forceDeviceID)
 {
        //Find best CUDA device, initialize and allocate memory
 
@@ -184,7 +201,7 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
 
                int deviceOK = fCudaDeviceProp.major < 9 && !(fCudaDeviceProp.major < 1 || (fCudaDeviceProp.major == 1 && fCudaDeviceProp.minor < 2)) && free >= fGPUMemSize;
 
-               if (fDebugLevel >= 2) HLTInfo("%s%2d: %s (Rev: %d.%d - Mem Avail %d / %d)%s", deviceOK ? " " : "[", i, fCudaDeviceProp.name, fCudaDeviceProp.major, fCudaDeviceProp.minor, free, fCudaDeviceProp.totalGlobalMem, deviceOK ? "" : " ]");
+               if (fDebugLevel >= 2) HLTInfo("%s%2d: %s (Rev: %d.%d - Mem Avail %d / %lld)%s", deviceOK ? " " : "[", i, fCudaDeviceProp.name, fCudaDeviceProp.major, fCudaDeviceProp.minor, free, (long long int) fCudaDeviceProp.totalGlobalMem, deviceOK ? "" : " ]");
                deviceSpeed = (long long int) fCudaDeviceProp.multiProcessorCount * (long long int) fCudaDeviceProp.clockRate * (long long int) fCudaDeviceProp.warpSize * (long long int) free;
                if (deviceOK && deviceSpeed > bestDeviceSpeed)
                {
@@ -214,19 +231,19 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
   if (fDebugLevel >= 1)
   {
          HLTInfo("Using CUDA Device %s with Properties:", fCudaDeviceProp.name);
-         HLTInfo("totalGlobalMem = %d", fCudaDeviceProp.totalGlobalMem);
-         HLTInfo("sharedMemPerBlock = %d", fCudaDeviceProp.sharedMemPerBlock);
+         HLTInfo("totalGlobalMem = %lld", (unsigned long long int) fCudaDeviceProp.totalGlobalMem);
+         HLTInfo("sharedMemPerBlock = %lld", (unsigned long long int) fCudaDeviceProp.sharedMemPerBlock);
          HLTInfo("regsPerBlock = %d", fCudaDeviceProp.regsPerBlock);
          HLTInfo("warpSize = %d", fCudaDeviceProp.warpSize);
-         HLTInfo("memPitch = %d", fCudaDeviceProp.memPitch);
+         HLTInfo("memPitch = %lld", (unsigned long long int) fCudaDeviceProp.memPitch);
          HLTInfo("maxThreadsPerBlock = %d", fCudaDeviceProp.maxThreadsPerBlock);
          HLTInfo("maxThreadsDim = %d %d %d", fCudaDeviceProp.maxThreadsDim[0], fCudaDeviceProp.maxThreadsDim[1], fCudaDeviceProp.maxThreadsDim[2]);
          HLTInfo("maxGridSize = %d %d %d", fCudaDeviceProp.maxGridSize[0], fCudaDeviceProp.maxGridSize[1], fCudaDeviceProp.maxGridSize[2]);
-         HLTInfo("totalConstMem = %d", fCudaDeviceProp.totalConstMem);
+         HLTInfo("totalConstMem = %lld", (unsigned long long int) fCudaDeviceProp.totalConstMem);
          HLTInfo("major = %d", fCudaDeviceProp.major);
          HLTInfo("minor = %d", fCudaDeviceProp.minor);
          HLTInfo("clockRate %d= ", fCudaDeviceProp.clockRate);
-         HLTInfo("textureAlignment %d= ", fCudaDeviceProp.textureAlignment);
+         HLTInfo("textureAlignment %lld= ", (unsigned long long int) fCudaDeviceProp.textureAlignment);
   }
 
   if (fCudaDeviceProp.major < 1 || (fCudaDeviceProp.major == 1 && fCudaDeviceProp.minor < 2))
@@ -298,7 +315,7 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
   }
 
   fCudaInitialized = 1;
-  HLTImportant("CUDA Initialisation successfull (Device %d: %s, Thread %dd)", cudaDevice, fCudaDeviceProp.name, fThreadId);
+  HLTImportant("CUDA Initialisation successfull (Device %d: %s, Thread %d)", cudaDevice, fCudaDeviceProp.name, fThreadId);
 
 #if defined(HLTCA_STANDALONE) & !defined(CUDA_DEVICE_EMULATION)
   if (fDebugLevel < 2)
@@ -351,7 +368,7 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
   return(0);
 }
 
-template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int alignment)
+template <class T> inline T* AliHLTTPCCAGPUTrackerNVCC::alignPointer(T* ptr, int alignment)
 {
        //Macro to align Pointers.
        //Will align to start at 1 MB segments, this should be consistent with every alignment in the tracker
@@ -365,7 +382,7 @@ template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int ali
        return((T*) adr);
 }
 
-bool AliHLTTPCCAGPUTracker::CudaFailedMsg(cudaError_t error)
+bool AliHLTTPCCAGPUTrackerNVCC::CudaFailedMsg(cudaError_t error)
 {
        //Check for CUDA Error and in the case of an error display the corresponding error string
        if (error == cudaSuccess) return(false);
@@ -373,7 +390,7 @@ bool AliHLTTPCCAGPUTracker::CudaFailedMsg(cudaError_t error)
        return(true);
 }
 
-int AliHLTTPCCAGPUTracker::CUDASync(char* state)
+int AliHLTTPCCAGPUTrackerNVCC::CUDASync(char* state)
 {
        //Wait for CUDA-Kernel to finish and check for CUDA errors afterwards
 
@@ -394,14 +411,14 @@ int AliHLTTPCCAGPUTracker::CUDASync(char* state)
        return(0);
 }
 
-void AliHLTTPCCAGPUTracker::SetDebugLevel(const int dwLevel, std::ostream* const NewOutFile)
+void AliHLTTPCCAGPUTrackerNVCC::SetDebugLevel(const int dwLevel, std::ostream* const NewOutFile)
 {
        //Set Debug Level and Debug output File if applicable
        fDebugLevel = dwLevel;
        if (NewOutFile) fOutFile = NewOutFile;
 }
 
-int AliHLTTPCCAGPUTracker::SetGPUTrackerOption(char* OptionName, int /*OptionValue*/)
+int AliHLTTPCCAGPUTrackerNVCC::SetGPUTrackerOption(char* OptionName, int /*OptionValue*/)
 {
        //Set a specific GPU Tracker Option
        {
@@ -413,7 +430,7 @@ int AliHLTTPCCAGPUTracker::SetGPUTrackerOption(char* OptionName, int /*OptionVal
 }
 
 #ifdef HLTCA_STANDALONE
-void AliHLTTPCCAGPUTracker::StandalonePerfTime(int iSlice, int i)
+void AliHLTTPCCAGPUTrackerNVCC::StandalonePerfTime(int iSlice, int i)
 {
   //Run Performance Query for timer i of slice iSlice
   if (fDebugLevel >= 1)
@@ -422,10 +439,10 @@ void AliHLTTPCCAGPUTracker::StandalonePerfTime(int iSlice, int i)
   }
 }
 #else
-void AliHLTTPCCAGPUTracker::StandalonePerfTime(int /*iSlice*/, int /*i*/) {}
+void AliHLTTPCCAGPUTrackerNVCC::StandalonePerfTime(int /*iSlice*/, int /*i*/) {}
 #endif
 
-void AliHLTTPCCAGPUTracker::DumpRowBlocks(AliHLTTPCCATracker* tracker, int iSlice, bool check)
+void AliHLTTPCCAGPUTrackerNVCC::DumpRowBlocks(AliHLTTPCCATracker* tracker, int iSlice, bool check)
 {
        //Dump Rowblocks to File
        if (fDebugLevel >= 4)
@@ -499,7 +516,7 @@ __global__ void PreInitRowBlocks(int4* const RowBlockPos, int* const RowBlockTra
                sliceDataHitWeights4[i] = i0;
 }
 
-int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTTPCCAClusterData* pClusterData, int firstSlice, int sliceCountLocal)
+int AliHLTTPCCAGPUTrackerNVCC::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTTPCCAClusterData* pClusterData, int firstSlice, int sliceCountLocal)
 {
        //Primary reconstruction function
        cudaStream_t* const cudaStreams = (cudaStream_t*) fpCudaStreams;
@@ -636,6 +653,14 @@ int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTT
                        AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&a);
 #endif
                fSlaveTrackers[firstSlice + iSlice].ReadEvent(&pClusterData[iSlice]);
+
+                 if (fDebugLevel >= 4)
+                 {
+                         *fOutFile << std::endl << std::endl << "Slice: " << fSlaveTrackers[firstSlice + iSlice].Param().ISlice() << std::endl;
+                         *fOutFile << "Slice Data:" << std::endl;
+                         fSlaveTrackers[firstSlice + iSlice].DumpSliceData(*fOutFile);
+                 }
+
 #ifdef HLTCA_GPU_TIME_PROFILE
                        AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&b);
                printf("Read %f %f\n", ((double) b - (double) a) / (double) c, ((double) a - (double) d) / (double) c);
@@ -938,7 +963,7 @@ RestartTrackletConstructor:
        return(0);
 }
 
-int AliHLTTPCCAGPUTracker::InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param)
+int AliHLTTPCCAGPUTrackerNVCC::InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param)
 {
        //Initialize Slice Tracker Parameter for a slave tracker
        fSlaveTrackers[iSlice].Initialize(param);
@@ -950,7 +975,7 @@ int AliHLTTPCCAGPUTracker::InitializeSliceParam(int iSlice, AliHLTTPCCAParam &pa
        return(0);
 }
 
-int AliHLTTPCCAGPUTracker::ExitGPU()
+int AliHLTTPCCAGPUTrackerNVCC::ExitGPU()
 {
        //Uninitialize CUDA
        cudaThreadSynchronize();
@@ -981,7 +1006,7 @@ int AliHLTTPCCAGPUTracker::ExitGPU()
        return(0);
 }
 
-void AliHLTTPCCAGPUTracker::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val)
+void AliHLTTPCCAGPUTrackerNVCC::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val)
 {
        //Set Output Control Pointers
        fOutputControl = val;
@@ -991,7 +1016,7 @@ void AliHLTTPCCAGPUTracker::SetOutputControl( AliHLTTPCCASliceOutput::outputCont
        }
 }
 
-int AliHLTTPCCAGPUTracker::GetThread()
+int AliHLTTPCCAGPUTrackerNVCC::GetThread()
 {
        //Get Thread ID
 #ifdef R__WIN32
@@ -1001,4 +1026,30 @@ int AliHLTTPCCAGPUTracker::GetThread()
 #endif
 }
 
-#endif
+unsigned long long int* AliHLTTPCCAGPUTrackerNVCC::PerfTimer(int iSlice, unsigned int i)
+{
+       //Returns pointer to PerfTimer i of slice iSlice
+       return(fSlaveTrackers ? fSlaveTrackers[iSlice].PerfTimer(i) : NULL);
+}
+
+const AliHLTTPCCASliceOutput::outputControlStruct* AliHLTTPCCAGPUTrackerNVCC::OutputControl() const
+{
+       //Return Pointer to Output Control Structure
+       return fOutputControl;
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::GetSliceCount() const
+{
+       //Return max slice count processable
+       return(fSliceCount);
+}
+
+AliHLTTPCCAGPUTracker* AliHLTTPCCAGPUTrackerNVCCCreate()
+{
+       return new AliHLTTPCCAGPUTrackerNVCC;
+}
+void AliHLTTPCCAGPUTrackerNVCCDestroy(AliHLTTPCCAGPUTracker* ptr)
+{
+       delete ptr;
+}
+
diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cxx
new file mode 100644 (file)
index 0000000..5fa74b7
--- /dev/null
@@ -0,0 +1,132 @@
+// **************************************************************************
+// This file is property of and copyright by the ALICE HLT Project          *
+// ALICE Experiment at CERN, All rights reserved.                           *
+//                                                                          *
+// Primary Authors: Sergey Gorbunov <sergey.gorbunov@kip.uni-heidelberg.de> *
+//                  Ivan Kisel <kisel@kip.uni-heidelberg.de>                *
+//                                     David Rohr <drohr@kip.uni-heidelberg.de>                                *
+//                  for The ALICE HLT Project.                              *
+//                                                                          *
+// Permission to use, copy, modify and distribute this software and its     *
+// documentation strictly for non-commercial purposes is hereby granted     *
+// without fee, provided that the above copyright notice appears in all     *
+// copies and that both the copyright notice and this permission notice     *
+// appear in the supporting documentation. The authors make no claims       *
+// about the suitability of this software for any purpose. It is            *
+// provided "as is" without express or implied warranty.                    *
+//                                                                          *
+//***************************************************************************
+
+#include "AliHLTTPCCAGPUTrackerNVCC.h"
+
+AliHLTTPCCAGPUTrackerNVCC::AliHLTTPCCAGPUTrackerNVCC() :
+       fGpuTracker(NULL),
+       fGPUMemory(NULL),
+       fHostLockedMemory(NULL),
+       fDebugLevel(0),
+       fOutFile(NULL),
+       fGPUMemSize(0),
+       fpCudaStreams(NULL),
+       fSliceCount(0),
+       fOutputControl(NULL),
+       fThreadId(0),
+       fCudaInitialized(0)
+       {};
+
+AliHLTTPCCAGPUTrackerNVCC::~AliHLTTPCCAGPUTrackerNVCC() {};
+
+void AliHLTTPCCAGPUTrackerNVCC::ReleaseGlobalLock(void* sem)
+{
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::CheckMemorySizes(int sliceCount)
+{
+  return(0);
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::InitGPU(int sliceCount, int forceDeviceID)
+{
+  return(0);
+}
+
+template <class T> inline T* AliHLTTPCCAGPUTrackerNVCC::alignPointer(T* ptr, int alignment)
+{
+       return((T*) NULL);
+}
+
+bool AliHLTTPCCAGPUTrackerNVCC::CudaFailedMsg(cudaError_t error)
+{
+       return(true);
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::CUDASync(char* state)
+{
+       return(0);
+}
+
+void AliHLTTPCCAGPUTrackerNVCC::SetDebugLevel(const int dwLevel, std::ostream* const NewOutFile)
+{
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::SetGPUTrackerOption(char* OptionName, int /*OptionValue*/)
+{
+       return(0);
+}
+
+void AliHLTTPCCAGPUTrackerNVCC::StandalonePerfTime(int /*iSlice*/, int /*i*/) {}
+
+void AliHLTTPCCAGPUTrackerNVCC::DumpRowBlocks(AliHLTTPCCATracker* tracker, int iSlice, bool check)
+{
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTTPCCAClusterData* pClusterData, int firstSlice, int sliceCountLocal)
+{
+       return(0);
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param)
+{
+       return(0);
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::ExitGPU()
+{
+       return(0);
+}
+
+void AliHLTTPCCAGPUTrackerNVCC::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val)
+{
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::GetThread()
+{
+    return(0);
+}
+
+unsigned long long int* AliHLTTPCCAGPUTrackerNVCC::PerfTimer(int iSlice, unsigned int i)
+{
+    static unsigned long long int tmp;
+    return(&tmp);
+}
+
+const AliHLTTPCCASliceOutput::outputControlStruct* AliHLTTPCCAGPUTrackerNVCC::OutputControl() const
+{
+       //Return Pointer to Output Control Structure
+       return fOutputControl;
+}
+
+int AliHLTTPCCAGPUTrackerNVCC::GetSliceCount() const
+{
+       //Return max slice count processable
+       return(fSliceCount);
+}
+
+AliHLTTPCCAGPUTracker* AliHLTTPCCAGPUTrackerNVCCCreate()
+{
+       return new AliHLTTPCCAGPUTrackerNVCC;
+} 
+void AliHLTTPCCAGPUTrackerNVCCDestroy(AliHLTTPCCAGPUTracker* ptr)
+{
+       delete ptr;
+}
+
diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.h
new file mode 100644 (file)
index 0000000..f9454a4
--- /dev/null
@@ -0,0 +1,98 @@
+// ************************************************************************
+// 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 ALIHLTTPCCAGPUTRACKERNVCC_H
+#define ALIHLTTPCCAGPUTRACKERNVCC_H
+
+#include "AliHLTTPCCAGPUTracker.h"
+#include "AliHLTTPCCADef.h"
+#include "AliHLTTPCCATracker.h"
+#include "AliHLTLogging.h"
+#include "AliHLTTPCCASliceOutput.h"
+
+class AliHLTTPCCARow;
+
+class AliHLTTPCCAGPUTrackerNVCC : public AliHLTTPCCAGPUTracker, public AliHLTLogging
+{
+public:
+       AliHLTTPCCAGPUTrackerNVCC();
+       virtual ~AliHLTTPCCAGPUTrackerNVCC();
+
+       virtual int InitGPU(int sliceCount = 12, int forceDeviceID = -1);
+       virtual int Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTTPCCAClusterData* pClusterData, int fFirstSlice, int fSliceCount = -1);
+       virtual int ExitGPU();
+
+       virtual void SetDebugLevel(const int dwLevel, std::ostream* const NewOutFile = NULL);
+       virtual int SetGPUTrackerOption(char* OptionName, int OptionValue);
+
+       virtual unsigned long long int* PerfTimer(int iSlice, unsigned int i);
+
+       virtual int InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param);
+       virtual void SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val);
+
+       virtual const AliHLTTPCCASliceOutput::outputControlStruct* OutputControl() const;
+       virtual int GetSliceCount() const;
+
+private:
+       static void* RowMemory(void* const BaseMemory, int iSlice) { return( ((char*) BaseMemory) + iSlice * sizeof(AliHLTTPCCARow) * (HLTCA_ROW_COUNT + 1) ); }
+       static void* CommonMemory(void* const BaseMemory, int iSlice) { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + iSlice * AliHLTTPCCATracker::CommonMemorySize() ); }
+       static void* SliceDataMemory(void* const BaseMemory, int iSlice) { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + iSlice * HLTCA_GPU_SLICE_DATA_MEMORY ); }
+       void* GlobalMemory(void* const BaseMemory, int iSlice) const { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + fSliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY) + iSlice * HLTCA_GPU_GLOBAL_MEMORY ); }
+       void* TracksMemory(void* const BaseMemory, int iSlice) const { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + fSliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY) + iSlice * HLTCA_GPU_TRACKS_MEMORY ); }
+       void* TrackerMemory(void* const BaseMemory, int iSlice) const { return( ((char*) BaseMemory) + HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + fSliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY + HLTCA_GPU_TRACKS_MEMORY) + iSlice * sizeof(AliHLTTPCCATracker) ); }
+
+       void DumpRowBlocks(AliHLTTPCCATracker* tracker, int iSlice, bool check = true);
+       int GetThread();
+       void ReleaseGlobalLock(void* sem);
+       int CheckMemorySizes(int sliceCount);
+
+       AliHLTTPCCATracker *fGpuTracker;
+       void* fGPUMemory;
+       void* fHostLockedMemory;
+
+       int CUDASync(char* state = "UNKNOWN");
+       template <class T> T* alignPointer(T* ptr, int alignment);
+
+       void StandalonePerfTime(int iSlice, int i);
+
+       int fDebugLevel;                        //Debug Level for GPU Tracker
+       std::ostream* fOutFile;         //Debug Output Stream Pointer
+       unsigned long long int fGPUMemSize;     //Memory Size to allocate on GPU
+
+       void* fpCudaStreams;
+
+       int fSliceCount;
+
+       static const int fgkNSlices = 36;
+       AliHLTTPCCATracker fSlaveTrackers[fgkNSlices];
+#ifdef HLTCA_GPUCODE
+       bool CudaFailedMsg(cudaError_t error);
+#endif //HLTCA_GPUCODE
+
+       AliHLTTPCCASliceOutput::outputControlStruct* fOutputControl;
+       
+       static bool fgGPUUsed;
+       int fThreadId;
+       int fCudaInitialized;
+
+       // disable copy
+       AliHLTTPCCAGPUTrackerNVCC( const AliHLTTPCCAGPUTrackerNVCC& );
+       AliHLTTPCCAGPUTrackerNVCC &operator=( const AliHLTTPCCAGPUTrackerNVCC& );
+
+       ClassDef( AliHLTTPCCAGPUTrackerNVCC, 0 )
+};
+
+#ifdef R__WIN32
+#define DLL_EXPORT __declspec(dllexport)
+#else
+#define DLL_EXPORT
+#endif
+
+extern "C" DLL_EXPORT AliHLTTPCCAGPUTracker* AliHLTTPCCAGPUTrackerNVCCCreate();
+extern "C" DLL_EXPORT void AliHLTTPCCAGPUTrackerNVCCDestroy(AliHLTTPCCAGPUTracker* ptr);
+
+#endif //ALIHLTTPCCAGPUTRACKER_H
index bb48864..e591ca1 100644 (file)
@@ -119,4 +119,4 @@ class AliHLTTPCCAGlobalMergerComponent : public AliHLTProcessor
     ClassDef( AliHLTTPCCAGlobalMergerComponent, 0 )
 };
 
-#endif
+#endif //ALIHLTTPCCAGLOBALMERGERCOMPONENT_H
index e08bf7d..056d815 100644 (file)
@@ -56,4 +56,4 @@ class AliHLTTPCCAGrid
 
 };
 
-#endif
+#endif //ALIHLTTPCCAGRID_H
index f010439..804608c 100644 (file)
@@ -35,4 +35,4 @@ class AliHLTTPCCAHit
 };
 
 
-#endif
+#endif //ALIHLTTPCCAHIT_H
index 4084a2f..920d3dd 100644 (file)
@@ -77,4 +77,4 @@ class AliHLTTPCCAHitArea
     int fHitOffset; // global hit offset XXX what's that?
 };
 
-#endif
+#endif //ALIHLTTPCCAHITAREA_H
index 4862d35..6dab028 100644 (file)
@@ -69,4 +69,4 @@ class AliHLTTPCCAMCPoint
     int   fTrackID;   //* mc track number
 };
 
-#endif
+#endif //ALIHLTTPCCAMCPOINT_H
index 14c326e..2f22927 100644 (file)
@@ -67,4 +67,4 @@ class AliHLTTPCCAMCTrack
 
 };
 
-#endif
+#endif //AliHLTTPCCAMCTrack
index 8d7f825..5b6bfa2 100644 (file)
@@ -15,7 +15,7 @@
 #include <math.h>
 #else
 #include "TMath.h"
-#endif
+#endif //HLTCA_STANDALONE | HLTCA_GPUCODE
 
 /**
  * @class ALIHLTTPCCAMath
@@ -65,7 +65,7 @@ typedef AliHLTTPCCAMath CAMath;
 #define choice(c1,c2,c3) c2
 #else
 #define choice(c1,c2,c3) c3
-#endif
+#endif //HLTCA_GPUCODE
 
 GPUd() inline float2 AliHLTTPCCAMath::MakeFloat2( float x, float y )
 {
@@ -74,7 +74,7 @@ GPUd() inline float2 AliHLTTPCCAMath::MakeFloat2( float x, float y )
   return ret;
 #else
   return make_float2( x, y );
-#endif
+#endif //HLTCA_GPUCODE
 }
 
 
@@ -92,7 +92,7 @@ GPUd() inline int AliHLTTPCCAMath::Nint( float x )
   return i;
 #else
   return TMath::Nint( x );
-#endif
+#endif //HLTCA_STANDALONE | HLTCA_GPUCODE
 }
 
 GPUd() inline bool AliHLTTPCCAMath::Finite( float x )
@@ -113,7 +113,7 @@ GPUd() inline float AliHLTTPCCAMath::Copysign( float x, float y )
 #else
   x = CAMath::Abs( x );
   return ( y >= 0 ) ? x : -x;
-#endif
+#endif //HLTCA_GPUCODE
 }
 
 
@@ -202,7 +202,7 @@ GPUd()  inline int AliHLTTPCCAMath::AtomicExch( int *addr, int val )
   int old = *addr;
   *addr = val;
   return old;
-#endif
+#endif //HLTCA_GPUCODE
 }
 
 GPUd()  inline int AliHLTTPCCAMath::AtomicAdd ( int *addr, int val )
@@ -213,7 +213,7 @@ GPUd()  inline int AliHLTTPCCAMath::AtomicAdd ( int *addr, int val )
   int old = *addr;
   *addr += val;
   return old;
-#endif
+#endif //HLTCA_GPUCODE
 }
 
 GPUd()  inline int AliHLTTPCCAMath::AtomicMax ( int *addr, int val )
@@ -224,7 +224,7 @@ GPUd()  inline int AliHLTTPCCAMath::AtomicMax ( int *addr, int val )
   int old = *addr;
   if ( *addr < val ) *addr = val;
   return old;
-#endif
+#endif //HLTCA_GPUCODE
 }
 
 GPUd()  inline int AliHLTTPCCAMath::AtomicMin ( int *addr, int val )
@@ -235,9 +235,9 @@ GPUd()  inline int AliHLTTPCCAMath::AtomicMin ( int *addr, int val )
   int old = *addr;
   if ( *addr > val ) *addr = val;
   return old;
-#endif
+#endif //HLTCA_GPUCODE
 }
 
 #undef CHOICE
 
-#endif
+#endif //ALIHLTTPCCAMATH_H
index 1c8b6d0..eb6ff37 100644 (file)
@@ -54,4 +54,4 @@ class AliHLTTPCCAMergedTrack
 };
 
 
-#endif
+#endif //ALIHLTTPCCAMERGEDTRACK_H
index 19ecbe1..39b683a 100644 (file)
@@ -15,7 +15,7 @@
 
 #if !defined(HLTCA_GPUCODE)
 #include <iostream>
-#endif
+#endif //HLTCA_GPUCODE
 
 class AliHLTTPCCASliceTrack;
 class AliHLTTPCCASliceOutput;
@@ -120,4 +120,4 @@ class AliHLTTPCCAMerger
     int fSliceNTrackInfos[fgkNSlices];                //* N of slice track infos in fTrackInfos array;
 };
 
-#endif
+#endif //ALIHLTTPCCAMERGER_H
index 763eed6..1bb3613 100644 (file)
@@ -88,4 +88,4 @@ GPUhd() inline void AliHLTTPCCAMergerOutput::SetPointers()
   fClusterPackedAmp  = ( UChar_t* ) ( fClusterId + fNTrackClusters );
 }
 
-#endif
+#endif //ALIHLTTPCCAMERGEROUTPUT_H
index ad5c1bc..e9ed78c 100644 (file)
@@ -31,7 +31,7 @@ class AliHLTTPCCANeighboursCleaner
         AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/ )
             : fIRow( 0 ), fIRowUp( 0 ), fIRowDn( 0 ), fNRows( 0 ), fNHits( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
-#endif
+#endif //!HLTCA_GPUCODE
 
       protected:
         int fIRow; // current row index
@@ -48,4 +48,4 @@ class AliHLTTPCCANeighboursCleaner
 };
 
 
-#endif
+#endif //ALIHLTTPCCANEIGHBOURSCLEANER_H
index a1bf697..45a743e 100644 (file)
@@ -27,7 +27,7 @@
 
 #ifdef DRAW
 #include "AliHLTTPCCADisplay.h"
-#endif
+#endif //DRAW
 
 GPUd() void AliHLTTPCCANeighboursFinder::Thread
 ( int /*nBlocks*/, int nThreads, int iBlock, int iThread, int iSync,
@@ -147,21 +147,12 @@ GPUd() void AliHLTTPCCANeighboursFinder::Thread
 
        for ( int ih = iThread; ih < s.fNHits; ih += nThreads ) {
 
-      unsigned short *neighUp = s.fB[iThread];
-      float2 *yzUp = s.fA[iThread];
-#if defined(HLTCA_GPUCODE) & kMaxN > ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP
-         unsigned short neighUp2[kMaxN - ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP];
-         float2 yzUp2[kMaxN - ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP];
-#endif
-      //unsigned short neighUp[5];
-      //float2 yzUp[5];
-
       int linkUp = -1;
       int linkDn = -1;
 
       if ( s.fDnNHits > 0 && s.fUpNHits > 0 ) {
 
-        int nNeighUp = 0;
+       
 
         // coordinates of the hit in the current row
 #if defined(HLTCA_GPU_TEXTURE_FETCHa)
@@ -173,6 +164,172 @@ GPUd() void AliHLTTPCCANeighboursFinder::Thread
         const float z = z0 + tracker.HitDataZ( row, ih ) * stepZ;
 #endif
 
+#ifdef FAST_NEIGHBOURS_FINDER
+
+//#define NFDEBUG
+
+#ifdef NFDEBUG
+       printf("\nSearching Neighbours for: %f %f\n", y, z);
+#endif
+
+       const float y0Up = rowUp.Grid().YMin();
+    const float z0Up = rowUp.Grid().ZMin();
+    const float stepYUp = rowUp.HstepY();
+    const float stepZUp = rowUp.HstepZ();
+
+    const float y0Dn = rowDn.Grid().YMin();
+    const float z0Dn = rowDn.Grid().ZMin();
+    const float stepYDn = rowDn.HstepY();
+    const float stepZDn = rowDn.HstepZ();
+
+               const float yMinUp = y*s.fUpTx - kAreaSize;
+               const float yMaxUp = y*s.fUpTx + kAreaSize;
+               const float yMinDn = y*s.fDnTx - kAreaSize;
+               const float yMaxDn = y*s.fDnTx + kAreaSize;
+               const float zMinUp = z*s.fUpTx - kAreaSize;
+               const float zMaxUp = z*s.fUpTx + kAreaSize;
+               const float zMinDn = z*s.fDnTx - kAreaSize;
+               const float zMaxDn = z*s.fDnTx + kAreaSize;
+
+               int bYUpMin, bZUpMin, bYDnMin, bZDnMin;
+               rowUp.Grid().GetBin(yMinUp, zMinUp, &bYUpMin, &bZUpMin);
+               rowDn.Grid().GetBin(yMinDn, zMinDn, &bYDnMin, &bZDnMin);
+
+               int bYUpMax, bZUpMax, bYDnMax, bZDnMax;
+               rowUp.Grid().GetBin(yMaxUp, zMaxUp, &bYUpMax, &bZUpMax);
+               rowDn.Grid().GetBin(yMaxDn, zMaxDn, &bYDnMax, &bZDnMax);
+
+               int nYUp = rowUp.Grid().Ny();
+               int nYDn = rowDn.Grid().Ny();
+
+               int ihUp = tracker.Data().FirstHitInBin(rowUp, bZUpMin * nYUp + bYUpMin);
+               int ihDn = bZDnMax * nYDn + bYDnMax >= rowDn.Grid().N() ? (rowDn.NHits() - 1) : (tracker.Data().FirstHitInBin(rowDn, bZDnMax * nYDn + bYDnMax + 1) - 1);
+
+               int ihUpMax = tracker.Data().FirstHitInBin(rowUp, bZUpMin * nYUp + bYUpMax + 1) - 1;
+               int ihDnMin = tracker.Data().FirstHitInBin(rowDn, bZDnMax * nYDn + bYDnMin);
+
+               float bestD = 1.e10;
+               int bestDn, bestUp;
+
+               int lastUp = 0, lastDn = 0;
+
+               while (true)
+               {
+                       float yUp = y0Up + tracker.HitDataY(rowUp, ihUp) * stepYUp;
+                       float zUp = z0Up + tracker.HitDataZ(rowUp, ihUp) * stepZUp;
+
+                       float yDn = y0Dn + tracker.HitDataY(rowDn, ihDn) * stepYDn;
+                       float zDn = z0Dn + tracker.HitDataZ(rowDn, ihDn) * stepZDn;
+
+                       int ihUpNext, ihDnNext;
+                       if (ihUp >= ihUpMax)
+                       {
+                               if (bZUpMin < bZUpMax)
+                               {
+                                       ihUpNext = tracker.Data().FirstHitInBin(rowUp, (bZUpMin + 1) * nYUp + bYUpMin);
+                               }
+                               else
+                               {
+                                       lastUp = 1;
+                               }
+                       }
+                       else
+                       {
+                               ihUpNext = ihUp + 1;
+                       }
+
+                       if (ihDn <= ihDnMin)
+                       {
+                               if (bZDnMax > bZDnMin)
+                               {
+                                       ihDnNext = tracker.Data().FirstHitInBin(rowDn, (bZDnMax - 1) * nYDn + bYDnMax);
+                               }
+                               else
+                               {
+                                       lastDn = 1;
+                               }
+                       }
+                       else
+                       {
+                               ihDnNext = ihDn - 1;
+                       }
+
+                       
+                       
+
+                       float dUp, dDn;
+                       if (!lastUp)
+                       {
+                               const float yUpNext = y0Up + tracker.HitDataY(rowUp, ihUpNext) * stepYUp;
+                               const float zUpNext = z0Up + tracker.HitDataZ(rowUp, ihUpNext) * stepZUp;
+                               const float dYUp = s.fUpDx * (yUpNext - y) - s.fDnDx * (yDn - y);
+                               const float dZUp = s.fUpDx * (zUpNext - y) - s.fDnDx * (zDn - z);
+#ifdef NFDEBUG
+                               printf("Checking Up y: %f nexty: %f z: %f nextz: %f\n", yUp, yUpNext, zUp, zUpNext);
+#endif
+                               dUp = dYUp * dYUp + dZUp * dZUp;
+                       }
+
+                       if (!lastDn)
+                       {
+                               const float yDnNext = y0Dn + tracker.HitDataY(rowDn, ihDnNext) * stepYDn;
+                               const float zDnNext = z0Dn + tracker.HitDataZ(rowDn, ihDnNext) * stepZDn;
+                               const float dYDn = s.fDnDx * (yDnNext - y) - s.fUpDx * (yUp - y);
+                               const float dZDn = s.fDnDx * (zDnNext - y) - s.fUpDx * (zUp - z);
+#ifdef NFDEBUG
+                               printf("Checking Dn y: %f nexty: %f z: %f nextz: %f\n", yDn, yDnNext, zDn, zDnNext);
+#endif
+                               dDn = dYDn * dYDn + dZDn * dZDn;
+                       }
+
+                       float d;
+                       if (lastDn || (dUp < dDn && !lastUp))
+                       {
+                               if (lastUp) break;
+                               d = dUp;
+                               if (ihUp >= ihUpMax)
+                               {
+                                       bZUpMin++;
+                                       ihUpMax = tracker.Data().FirstHitInBin(rowUp, bZUpMin * nYUp + bYUpMax + 1) - 1;
+                               }
+                               ihUp = ihUpNext;
+                       }
+                       else
+                       {
+                               d = dDn;
+                               if (ihDn <= ihDnMin)
+                               {
+                                       bZDnMax--;
+                                       ihDnMin = tracker.Data().FirstHitInBin(rowDn, bZDnMax * nYDn + bYDnMin);
+                               }
+                               ihDn = ihDnNext;
+                       }
+                       if (d < bestD)
+                       {
+                               bestD = d;
+                               bestUp = ihUp;
+                               bestDn = ihDn;
+                       }
+               }
+
+               if (bestD < chi2Cut)
+               {
+                       linkUp = bestUp;
+                       linkDn = bestDn;
+               }
+
+#else
+               //Old Slow Neighbours finder
+      unsigned short *neighUp = s.fB[iThread];
+      float2 *yzUp = s.fA[iThread];
+#if defined(HLTCA_GPUCODE) & kMaxN > ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP
+         unsigned short neighUp2[kMaxN - ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP];
+         float2 yzUp2[kMaxN - ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP];
+#endif
+      //unsigned short neighUp[5];
+      //float2 yzUp[5];
+
+               int nNeighUp = 0;
         AliHLTTPCCAHitArea areaDn, areaUp;
         areaUp.Init( rowUp, tracker.Data(), y*s.fUpTx, z*s.fUpTx, kAreaSize, kAreaSize );
         areaDn.Init( rowDn, tracker.Data(), y*s.fDnTx, z*s.fDnTx, kAreaSize, kAreaSize );
@@ -240,8 +397,8 @@ GPUd() void AliHLTTPCCANeighboursFinder::Thread
         }
 #ifdef DRAW
         std::cout << "n NeighUp = " << nNeighUp << ", n NeighDn = " << nNeighDn << std::endl;
-#endif
-
+#endif //DRAW
+#endif //FAST_NEIGHBOURS_FINDER
       }
 
       tracker.SetHitLinkUpData( row, ih, linkUp );
@@ -253,7 +410,7 @@ GPUd() void AliHLTTPCCANeighboursFinder::Thread
         AliHLTTPCCADisplay::Instance().DrawSliceHit( s.fIRow, ih, kBlue, 1. );
         AliHLTTPCCADisplay::Instance().Ask();
       }
-#endif
+#endif //DRAW
     }
   }
 }
index ff6dbee..6ea7726 100644 (file)
@@ -32,7 +32,7 @@ class AliHLTTPCCANeighboursFinder
         AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/ )
             : fNHits( 0 ), fUpNHits( 0 ), fDnNHits( 0 ), fUpDx( 0 ), fDnDx( 0 ), fUpTx( 0 ), fDnTx( 0 ), fIRow( 0 ), fIRowUp( 0 ), fIRowDn( 0 ), fNRows( 0 ), fRow(), fRowUp(), fRowDown() {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
-#endif
+#endif //!HLTCA_GPUCODE
 
       protected:
         int fNHits; // n hits
@@ -59,4 +59,4 @@ class AliHLTTPCCANeighboursFinder
 };
 
 
-#endif
+#endif //ALIHLTTPCCANEIGHBOURSFINDER_H
index 0c1553c..5cf8461 100644 (file)
@@ -58,4 +58,4 @@ class AliHLTTPCCAOutTrack
 };
 
 
-#endif
+#endif //ALIHLTTPCCAOUTTRACK_H
index 7a1c463..d3a7e97 100644 (file)
@@ -37,7 +37,7 @@ class AliHLTTPCCAParam
                             float rMin, float rMax, float zMin, float zMax,
                             float padPitch, float zSigma, float bz );
     void Update();
-#endif
+#endif //!HLTCA_GPUCODE
 
        GPUd() void Slice2Global( float x, float y,  float z,
                               float *X, float *Y,  float *Z ) const;
@@ -168,4 +168,4 @@ GPUd() inline float AliHLTTPCCAParam::GetBz( const AliHLTTPCCATrackParam &t ) co
   return GetBz( t.X(), t.Y(), t.Z() );
 }
 
-#endif
+#endif //ALIHLTTPCCAPARAM_H
index 8b0c245..872066f 100644 (file)
@@ -64,7 +64,7 @@ GPUg() void AliHLTTPCCAProcess( int nBlocks, int nThreads, AliHLTTPCCATracker &t
   }
 }
 
-#endif
+#endif //HLTCA_GPUCODE
 
 
 
@@ -105,6 +105,6 @@ GPUg() void AliHLTTPCCAProcess1( int nBlocks, int nThreads, AliHLTTPCCATracker &
   }
 }
 
-#endif
+#endif //HLTCA_GPUCODE
 
-#endif
+#endif //ALIHLTTPCCAPROCESS_H
index d04a579..b63062a 100644 (file)
@@ -27,7 +27,7 @@ class AliHLTTPCCARow
 
 #if !defined(HLTCA_GPUCODE)
     AliHLTTPCCARow();
-#endif
+#endif //!HLTCA_GPUCODE
 
     GPUhd() int   NHits()    const { return fNHits; }
     GPUhd() float X()        const { return fX; }
@@ -64,4 +64,4 @@ class AliHLTTPCCARow
     unsigned int fFirstHitInBinOffset; // offset in Tracker::fRowData to find the FirstHitInBin
 };
 
-#endif
+#endif //ALIHLTTPCCAROW_H
index 59e4bbc..2fefe60 100644 (file)
@@ -53,7 +53,7 @@ class AliHLTTPCCASliceData
 
 #ifndef HLTCA_GPUCODE
        ~AliHLTTPCCASliceData();
-#endif
+#endif //!HLTCA_GPUCODE
 
     void InitializeRows( const AliHLTTPCCAParam &parameters );
 
@@ -283,4 +283,4 @@ GPUd() inline int_v AliHLTTPCCASliceData::HitWeight( const AliHLTTPCCARow &row,
 
 typedef AliHLTTPCCASliceData SliceData;
 
-#endif // SLICEDATA_H
+#endif // ALIHLTTPCCASLICEDATA_H
index 0834d5d..a20b123 100644 (file)
@@ -46,4 +46,4 @@ class AliHLTTPCCASliceTrack
 };
 
 
-#endif
+#endif //ALIHLTTPCCASLICETRACK_H
index 0cdfa3e..1597ca0 100644 (file)
@@ -91,7 +91,7 @@ class AliHLTTPCCAStandaloneFramework
 #ifdef HLTCA_STANDALONE
        static inline void StandaloneQueryTime(unsigned long long int *i);
        static inline void StandaloneQueryFreq(unsigned long long int *i);
-#endif
+#endif //HLTCA_STANDALONE
 
   private:
 
@@ -123,7 +123,7 @@ class AliHLTTPCCAStandaloneFramework
                  timespec t;
                  clock_gettime(CLOCK_REALTIME, &t);
                  *i = (unsigned long long int) t.tv_sec * (unsigned long long int) 1000000000 + (unsigned long long int) t.tv_nsec;
-       #endif
+       #endif //R__WIN32
        }
 
        void AliHLTTPCCAStandaloneFramework::StandaloneQueryFreq(unsigned long long int *i)
@@ -132,8 +132,8 @@ class AliHLTTPCCAStandaloneFramework
                  QueryPerformanceFrequency((LARGE_INTEGER*) i);
        #else
                *i = 1000000000;
-       #endif
+       #endif //R__WIN32
        }
-#endif
+#endif //HLTCA_STANDALONE
 
-#endif
+#endif //ALIHLTTPCCASTANDALONEFRAMEWORK_H
index 1facf3c..3b6c48d 100644 (file)
@@ -32,7 +32,7 @@ class AliHLTTPCCAStartHitsFinder
         AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/ )
             : fIRow( 0 ), fNRows( 0 ), fNHits( 0 ), fNOldStartHits( 0 ), fNRowStartHits( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
-#endif
+#endif //!HLTCA_GPUCODE
 
       protected:
         int fIRow; // row index
@@ -50,4 +50,4 @@ class AliHLTTPCCAStartHitsFinder
 };
 
 
-#endif
+#endif //ALIHLTTPCCASTARTHITSFINDER_H
index 0969c26..a919d92 100644 (file)
@@ -32,7 +32,7 @@ class AliHLTTPCCAStartHitsSorter
         AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/ )
             : fStartRow( 0 ), fNRows( 0 ), fStartOffset( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
-#endif
+#endif //!HLTCA_GPUCODE
 
       protected:
         int fStartRow;         // start row index
@@ -47,4 +47,4 @@ class AliHLTTPCCAStartHitsSorter
 };
 
 
-#endif
+#endif //ALIHLTTPCCASTARTHITSSORTER_H
index f2e610d..2faefd7 100644 (file)
@@ -26,7 +26,7 @@ class AliHLTTPCCATrack
 #if !defined(HLTCA_GPUCODE)
     AliHLTTPCCATrack() : fAlive( 0 ), fFirstHitID( 0 ), fNHits( 0 ), fParam() {}
     ~AliHLTTPCCATrack() {}
-#endif
+#endif //!HLTCA_GPUCODE
 
     GPUhd() bool Alive()               const { return fAlive; }
     GPUhd() int  NHits()               const { return fNHits; }
@@ -50,4 +50,4 @@ class AliHLTTPCCATrack
     //ClassDef(AliHLTTPCCATrack,1)
 };
 
-#endif
+#endif //ALIHLTTPCCATRACK_H
index 365cf2c..df454a5 100644 (file)
@@ -84,4 +84,4 @@ GPUd() inline void AliHLTTPCCATrackLinearisation::Set( float SinPhi1, float CosP
   SetQPt( QPt1 );
 }
 
-#endif
+#endif //ALIHLTTPCCATRACKLINEARISATION_H
index 59a1bda..c2c2c11 100644 (file)
@@ -139,7 +139,7 @@ class AliHLTTPCCATrackParam
 
 #ifndef HLTCA_GPUCODE
   private:
-#endif
+#endif //!HLTCA_GPUCODE
        AliHLTTPCCATrackParam2 fParam; // Track Parameters
 
   private:
@@ -178,4 +178,4 @@ GPUd() inline void AliHLTTPCCATrackParam::InitParam()
   SetCov( 14, 10. );
 }
 
-#endif
+#endif //ALIHLTTPCCATRACKPARAM_H
index 57b908b..e2a37e2 100755 (executable)
@@ -63,4 +63,4 @@ class AliHLTTPCCATrackParam2
     float fP[5];   // 'active' track parameters: Y, Z, SinPhi, DzDs, q/Pt
 };
 
-#endif
+#endif //ALIHLTTPCCATRACKPARAM2_H
index 2d6b0aa..fc24d4a 100644 (file)
@@ -51,7 +51,7 @@
 
 #ifdef DRAW1
 #include "AliHLTTPCCADisplay.h"
-#endif //DRAW
+#endif //DRAW1
 
 #ifdef HLTCA_INTERNAL_PERFORMANCE
 //#include "AliHLTTPCCAPerformance.h"
@@ -511,18 +511,18 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
     AliHLTTPCCADisplay::Instance().Ask();
   }
   }
-#endif
+#endif //DRAW1
 
        fCommonMem->fNTracklets = fCommonMem->fNTracks = fCommonMem->fNTrackHits = 0;
 
 #if !defined(HLTCA_GPUCODE)
 
-  /*if (fGPUDebugLevel >= 6)
+  if (fGPUDebugLevel >= 6)
   {
-         *fGPUDebugOut << endl << endl << "Slice: " << Param().ISlice() << endl;
-         *fGPUDebugOut << "Slice Data:" << endl;
+         *fGPUDebugOut << std::endl << std::endl << "Slice: " << Param().ISlice() << std::endl;
+         *fGPUDebugOut << "Slice Data:" << std::endl;
          DumpSliceData(*fGPUDebugOut);
-  }*/
+  }
 
   StandalonePerfTime(1);
 
@@ -543,7 +543,7 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
     AliHLTTPCCADisplay::Instance().DrawSliceLinks( -1, -1, 1 );
     AliHLTTPCCADisplay::Instance().Ask();
   }
-#endif
+#endif //DRAW1
 
   RunNeighboursCleaner();
 
@@ -621,7 +621,7 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
     }
     disp.Ask();
   }
-#endif
+#endif //DRAW1
 
   timer0.Stop();
   fTimers[0] = timer0.CpuTime() / 100.;
index cf675f1..293b46a 100644 (file)
@@ -111,7 +111,7 @@ class AliHLTTPCCATracker
   
 #if !defined(HLTCA_GPUCODE)
   void Reconstruct();
-#endif
+#endif //!HLTCA_GPUCODE
   
   //Make Reconstruction steps directly callable (Used for GPU debugging)
   void RunNeighboursFinder();
@@ -157,7 +157,7 @@ class AliHLTTPCCATracker
   GPUh() void WriteEvent( std::ostream &out );
   GPUh() void WriteTracks( std::ostream &out ) ;
   GPUh() void ReadTracks( std::istream &in );
-#endif
+#endif //!HLTCA_GPUCODE
   
   GPUhd() const AliHLTTPCCAParam &Param() const { return fParam; }
   GPUhd() void SetParam( const AliHLTTPCCAParam &v ) { fParam = v; }
@@ -264,7 +264,7 @@ private:
   
 #ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
   char* fStageAtSync;                          //Pointer to array storing current stage for every thread at every sync point
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
   
   AliHLTTPCCAParam fParam; // parameters
   double fTimers[10]; // timers
@@ -325,4 +325,4 @@ private:
 };
 
 
-#endif
+#endif //ALIHLTTPCCATRACKER_H
index 536fab6..24577bb 100644 (file)
@@ -236,7 +236,7 @@ int AliHLTTPCCATrackerComponent::ReadConfigurationString(  const char* arguments
       HLTInfo( "The special output type \"TRAKSEGS\" is set" );
       continue;
     }
-    
+
     if (argument.CompareTo( "-allowGPU" ) == 0) {
       fAllowGPU = 1;
       HLTImportant( "Will try to run tracker on GPU" );
index 82e96ea..1b05fe0 100644 (file)
@@ -82,7 +82,7 @@ class AliHLTTPCCATrackerComponent : public AliHLTProcessor
     double fSolenoidBz;                                            // see above
     int fMinNTrackClusters; //* required min number of clusters on the track
     double fClusterZCut;  //* cut on cluster Z position (for noise rejection at the age of TPC)
-  double fNeighboursSearchArea; //* area in cm for the neighbour search algorithm
+       double fNeighboursSearchArea; //* area in cm for the neighbour search algorithm
     double fClusterErrorCorrectionY; // correction for the cluster errors
     double fClusterErrorCorrectionZ; // correction for the cluster errors
 
@@ -90,7 +90,7 @@ class AliHLTTPCCATrackerComponent : public AliHLTProcessor
     double fRecoTime; //* total reconstruction time [s]
     Long_t    fNEvents;  //* number of reconstructed events
     bool fOutputTRAKSEGS; //* use old type of output
-    bool fAllowGPU;    //* Allow this tracker to run on GPU
+    bool fAllowGPU;    //* Allow this tracker to run on GPU
 
     static bool CompareClusters( AliHLTTPCSpacePointData *a, AliHLTTPCSpacePointData *b );
 
@@ -103,4 +103,4 @@ class AliHLTTPCCATrackerComponent : public AliHLTProcessor
     ClassDef( AliHLTTPCCATrackerComponent, 0 );
 
 };
-#endif
+#endif //ALIHLTTPCCATRACKERCOMPONENT_H
index 3e881a7..ac2748d 100644 (file)
 #include "AliHLTTPCCAMath.h"
 #include "AliHLTTPCCAClusterData.h"
 
+#ifdef R__WIN32
+#include <windows.h>
+#include <winbase.h>
+#else
+#include <dlfcn.h>
+#endif
+
 #ifdef HLTCA_STANDALONE
 #include <omp.h>
 #endif
 
+ClassImp( AliHLTTPCCATrackerFramework )
+
 int AliHLTTPCCATrackerFramework::InitGPU(int sliceCount, int forceDeviceID)
 {
        //Initialize GPU Tracker and determine if GPU available
        int retVal;
+       if (!fGPULibAvailable)
+       {
+               HLTError("GPU Library not loaded\n");
+               return(1);
+       }
        if (fGPUTrackerAvailable && (retVal = ExitGPU())) return(retVal);
-       retVal = fGPUTracker.InitGPU(sliceCount, forceDeviceID);
+       retVal = fGPUTracker->InitGPU(sliceCount, forceDeviceID);
        fUseGPUTracker = fGPUTrackerAvailable = retVal == 0;
        fGPUSliceCount = sliceCount;
        return(retVal);
@@ -46,13 +60,13 @@ int AliHLTTPCCATrackerFramework::ExitGPU()
        if (!fGPUTrackerAvailable) return(0);
        fUseGPUTracker = false;
        fGPUTrackerAvailable = false;
-       return(fGPUTracker.ExitGPU());
+       return(fGPUTracker->ExitGPU());
 }
 
 void AliHLTTPCCATrackerFramework::SetGPUDebugLevel(int Level, std::ostream *OutFile, std::ostream *GPUOutFile)
 {
        //Set Debug Level for GPU Tracker and also for CPU Tracker for comparison reasons
-       fGPUTracker.SetDebugLevel(Level, GPUOutFile);
+       fGPUTracker->SetDebugLevel(Level, GPUOutFile);
        fGPUDebugLevel = Level;
        for (int i = 0;i < fgkNSlices;i++)
        {
@@ -76,7 +90,7 @@ GPUhd() void AliHLTTPCCATrackerFramework::SetOutputControl( AliHLTTPCCASliceOutp
 {
        //Set Output Control Pointers
        fOutputControl = val;
-       fGPUTracker.SetOutputControl(val);
+       fGPUTracker->SetOutputControl(val);
        for (int i = 0;i < fgkNSlices;i++)
        {
                fCPUTrackers[i].SetOutputControl(val);
@@ -88,14 +102,14 @@ int AliHLTTPCCATrackerFramework::ProcessSlices(int firstSlice, int sliceCount, A
        //Process sliceCount slices starting from firstslice, in is pClusterData array, out pOutput array
        if (fUseGPUTracker)
        {
-               if (fGPUTracker.Reconstruct(pOutput, pClusterData, firstSlice, CAMath::Min(sliceCount, fgkNSlices - firstSlice))) return(1);
+               if (fGPUTracker->Reconstruct(pOutput, pClusterData, firstSlice, CAMath::Min(sliceCount, fgkNSlices - firstSlice))) return(1);
        }
        else
        {
 #ifdef HLTCA_STANDALONE
                if (fOutputControl->fOutputPtr && omp_get_max_threads() > 1)
                {
-                       printf("fOutputPtr must not be used with OpenMP\n");
+                       HLTError("fOutputPtr must not be used with OpenMP\n");
                        return(1);
                }
 
@@ -124,13 +138,114 @@ int AliHLTTPCCATrackerFramework::ProcessSlices(int firstSlice, int sliceCount, A
 unsigned long long int* AliHLTTPCCATrackerFramework::PerfTimer(int GPU, int iSlice, int iTimer)
 {
        //Performance information for slice trackers
-       return(GPU ? fGPUTracker.PerfTimer(iSlice, iTimer) : fCPUTrackers[iSlice].PerfTimer(iTimer));
+       return(GPU ? fGPUTracker->PerfTimer(iSlice, iTimer) : fCPUTrackers[iSlice].PerfTimer(iTimer));
 }
 
 int AliHLTTPCCATrackerFramework::InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param)
 {
        //Initialize Tracker Parameters for a slice
-       if (fGPUTrackerAvailable && fGPUTracker.InitializeSliceParam(iSlice, param)) return(1);
+       if (fGPUTrackerAvailable && fGPUTracker->InitializeSliceParam(iSlice, param)) return(1);
        fCPUTrackers[iSlice].Initialize(param);
        return(0);
 }
+
+#define GPULIBNAME "libAliHLTTPCCAGPU"
+
+AliHLTTPCCATrackerFramework::AliHLTTPCCATrackerFramework(int allowGPU) :       fGPULibAvailable(false), fGPUTrackerAvailable(false), fUseGPUTracker(false), fGPUDebugLevel(0), fGPUSliceCount(0), fGPUTracker(NULL), fGPULib(NULL), fOutputControl( NULL ), fCPUSliceCount(fgkNSlices)
+{
+       //Constructor
+#ifdef R__WIN32
+       HMODULE hGPULib = LoadLibraryEx(GPULIBNAME ".dll", NULL, NULL);
+#else
+       void* hGPULib = dlopen(GPULIBNAME ".so", RTLD_NOW);
+#endif
+       if (hGPULib == NULL)
+       {
+#ifndef R__WIN32
+               HLTInfo("%s\n", dlerror());
+#endif
+               if (allowGPU)
+               {
+                       HLTError("Error Opening cagpu library for GPU Tracker, will fallback to CPU\n");
+               }
+               else
+               {
+                       HLTInfo("Cagpu library was not found, Tracking on GPU will not be available");
+               }
+               fGPUTracker = new AliHLTTPCCAGPUTracker;
+       }
+       else
+       {
+#ifdef R__WIN32
+               FARPROC createFunc = GetProcAddress(hGPULib, "AliHLTTPCCAGPUTrackerNVCCCreate");
+#else
+               void* createFunc = (void*) dlsym(hGPULib, "AliHLTTPCCAGPUTrackerNVCCCreate");
+#endif
+               if (createFunc == NULL)
+               {
+                       HLTError("Error Creating GPU Tracker\n");
+#ifdef R__WIN32
+                       FreeLibrary(hGPULib);
+#else
+                       dlclose(hGPULib);
+#endif
+                       fGPUTracker = new AliHLTTPCCAGPUTracker;
+               }
+               else
+               {
+                       AliHLTTPCCAGPUTracker* (*tmp)() = (AliHLTTPCCAGPUTracker* (*)()) createFunc;
+                       fGPUTracker = tmp();
+                       fGPULibAvailable = true;
+                       fGPULib = (void*) (size_t) hGPULib;
+                       HLTImportant("GPU Tracker Created by Wrapper library");
+               }
+       }
+
+       if (allowGPU && fGPULibAvailable)
+       {
+               fUseGPUTracker = (fGPUTrackerAvailable= (fGPUTracker->InitGPU() == 0));
+               fGPUSliceCount = fGPUTrackerAvailable ? fGPUTracker->GetSliceCount() : 0;
+               HLTInfo("GPU Tracker Initialized and available in framework");
+       }
+}
+
+AliHLTTPCCATrackerFramework::~AliHLTTPCCATrackerFramework()
+{
+#ifdef R__WIN32
+       HMODULE hGPULib = (HMODULE) (size_t) fGPULib;
+#else
+       void* hGPULib = fGPULib;
+#endif
+       if (fGPULib)
+       {
+               if (fGPUTracker)
+               {
+#ifdef R__WIN32
+                       FARPROC destroyFunc = GetProcAddress(hGPULib, "AliHLTTPCCAGPUTrackerNVCCDestroy");
+#else
+                       void* destroyFunc = (void*) dlsym(hGPULib, "AliHLTTPCCAGPUTrackerNVCCDestroy");
+#endif
+                       if (destroyFunc == NULL)
+                       {
+                               HLTError("Error Freeing GPU Tracker\n");
+                       }
+                       else
+                       {
+                               void (*tmp)(AliHLTTPCCAGPUTracker*) =  (void (*)(AliHLTTPCCAGPUTracker*)) destroyFunc;
+                               tmp(fGPUTracker);
+                       }
+               }
+
+#ifdef R__WIN32
+               FreeLibrary(hGPULib);
+#else
+               dlclose(hGPULib);
+#endif
+       }
+       else if (fGPUTracker)
+       {
+               delete fGPUTracker;
+       }
+       fGPULib = NULL;
+       fGPUTracker = NULL;
+}
index acd7262..f9fc763 100644 (file)
 #include "AliHLTTPCCAGPUTracker.h"
 #include "AliHLTTPCCAParam.h"
 #include "AliHLTTPCCASliceOutput.h"
+#include "AliHLTLogging.h"
 #include <iostream>
 
 class AliHLTTPCCASliceOutput;
 class AliHLTTPCCAClusterData;
 
-class AliHLTTPCCATrackerFramework
+class AliHLTTPCCATrackerFramework : AliHLTLogging
 {
 public:
-       AliHLTTPCCATrackerFramework(int autoTryGPU = 1) :
-         fGPUTrackerAvailable(false), fUseGPUTracker(false), fGPUDebugLevel(0), fGPUSliceCount(0), fGPUTracker(), fOutputControl( NULL ), fCPUSliceCount(fgkNSlices)
-         {
-               if (autoTryGPU)
-               {
-                 fUseGPUTracker = (fGPUTrackerAvailable= (fGPUTracker.InitGPU() == 0));
-                 fGPUSliceCount = fGPUTrackerAvailable ? fGPUTracker.GetSliceCount() : 0;
-               }
-         }
-       ~AliHLTTPCCATrackerFramework()
-         {}
+       AliHLTTPCCATrackerFramework(int allowGPU = 1);
+       ~AliHLTTPCCATrackerFramework();
 
        int InitGPU(int sliceCount = 1, int forceDeviceID = -1);
        int ExitGPU();
        void SetGPUDebugLevel(int Level, std::ostream *OutFile = NULL, std::ostream *GPUOutFile = NULL);
-       int SetGPUTrackerOption(char* OptionName, int OptionValue) {return(fGPUTracker.SetGPUTrackerOption(OptionName, OptionValue));}
+       int SetGPUTrackerOption(char* OptionName, int OptionValue) {return(fGPUTracker->SetGPUTrackerOption(OptionName, OptionValue));}
        int SetGPUTracker(bool enable);
 
        int InitializeSliceParam(int iSlice, AliHLTTPCCAParam &param);
@@ -57,11 +49,13 @@ public:
 private:
   static const int fgkNSlices = 36;       //* N slices
 
+  bool fGPULibAvailable;       //Is the Library with the GPU code available at all?
   bool fGPUTrackerAvailable; // Is the GPU Tracker Available?
   bool fUseGPUTracker; // use the GPU tracker 
   int fGPUDebugLevel;  // debug level for the GPU code
   int fGPUSliceCount;  //How many slices to process parallel
-  AliHLTTPCCAGPUTracker fGPUTracker;
+  AliHLTTPCCAGPUTracker* fGPUTracker;
+  void* fGPULib;
 
   AliHLTTPCCASliceOutput::outputControlStruct* fOutputControl;
 
@@ -71,6 +65,8 @@ private:
   AliHLTTPCCATrackerFramework( const AliHLTTPCCATrackerFramework& );
   AliHLTTPCCATrackerFramework &operator=( const AliHLTTPCCATrackerFramework& );
 
+  ClassDef( AliHLTTPCCATrackerFramework, 0 )
+
 };
 
-#endif
+#endif //ALIHLTTPCCATRACKERFRAMEWORK_H
index 4dc99b6..d0cbab6 100644 (file)
@@ -27,7 +27,7 @@ class AliHLTTPCCATracklet
     AliHLTTPCCATracklet() : fNHits( 0 ), fFirstRow( 0 ), fLastRow( 0 ), fParam() {};
     void Dummy() const ;
     ~AliHLTTPCCATracklet() {}
-#endif
+#endif //!HLTCA_GPUCODE
 
     GPUhd() int  NHits()                const { return fNHits;      }
     GPUhd() int  FirstRow()             const { return fFirstRow;   }
@@ -36,15 +36,13 @@ class AliHLTTPCCATracklet
 #ifndef EXTERN_ROW_HITS
     GPUhd() int  RowHit( int i )   const { return fRowHits[i];    }
        GPUhd() const int* RowHits()    const                   { return(fRowHits); }
-#endif
+    GPUhd() void SetRowHit( int irow, int ih )  { fRowHits[irow] = ih;    }
+#endif //EXTERN_ROW_HITS
 
     GPUhd() void SetNHits( int v )               {  fNHits = v;      }
     GPUhd() void SetFirstRow( int v )            {  fFirstRow = v;   }
     GPUhd() void SetLastRow( int v )             {  fLastRow = v;    }
     GPUhd() void SetParam( const AliHLTTPCCATrackParam2 &v ) { fParam = v;      }
-#ifndef EXTERN_ROW_HITS
-    GPUhd() void SetRowHit( int irow, int ih )  { fRowHits[irow] = ih;    }
-#endif
 
   private:
     int fNHits;                 // N hits
@@ -53,7 +51,7 @@ class AliHLTTPCCATracklet
     AliHLTTPCCATrackParam2 fParam; // tracklet parameters
 #ifndef EXTERN_ROW_HITS
     int fRowHits[HLTCA_ROW_COUNT + 1];          // hit index for each TPC row
-#endif
+#endif //EXTERN_ROW_HITS
 };
 
-#endif
+#endif //ALIHLTTPCCATRACKLET_H
index 553d4df..d94f476 100644 (file)
@@ -34,7 +34,7 @@
 
 #ifdef DRAW
 #include "AliHLTTPCCADisplay.h"
-#endif
+#endif //DRAW
 
 #define kMaxRowGap 4
 
@@ -78,7 +78,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::ReadData
     for ( int i = iThread; i < numberOfHitsAligned; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
       sMem1a[i] = tracker.HitDataZ( row, i );
     }
-#endif
+#endif //HLTCA_GPU_REORDERHITDATA
 
     short *sMem2 = reinterpret_cast<short *>( s.fData[jr] ) + 2 * numberOfHitsAligned;
     for ( int i = iThread; i < numberOfHitsAligned; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
@@ -136,7 +136,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::ReadData
        {
                sharedMem[i] = sourceMem[i];
        }
-#endif
+#endif //!HLTCA_GPU_PREFETCHDATA
 }
 
 
@@ -147,7 +147,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
   &s
 #else
   &/*s*/
-#endif  
+#endif  //!HLTCA_GPUCODE
   , AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam )
 {
   // reconstruction of tracklets, tracklet store step
@@ -201,7 +201,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
         AliHLTTPCCADisplay::Instance().Ask();
       }
     }
-#endif
+#endif //DRAW
     if ( CAMath::Abs( tParam.Par()[4] ) < 1.e-4 ) tParam.SetPar( 4, 1.e-4 );
        if (r.fStartRow < r.fFirstRow) r.fFirstRow = r.fStartRow;
        tracklet.SetFirstRow( r.fFirstRow );
@@ -210,20 +210,20 @@ GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
     tracklet.SetParam( tParam.fParam );
 #else
     tracklet.SetParam( tParam.GetParam() );
-#endif
+#endif //HLTCA_GPUCODE
     int w = ( r.fNHits << 16 ) + r.fItr;
     for ( int iRow = r.fFirstRow; iRow <= r.fLastRow; iRow++ ) {
 #ifdef EXTERN_ROW_HITS
       int ih = tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr];
 #else
          int ih = tracklet.RowHit( iRow );
-#endif
+#endif //EXTERN_ROW_HITS
       if ( ih >= 0 ) {
 #if defined(HLTCA_GPUCODE) & !defined(HLTCA_GPU_PREFETCHDATA) & !defined(HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY)
            tracker.MaximizeHitWeight( s.fRows[ iRow ], ih, w );
 #else
            tracker.MaximizeHitWeight( tracker.Row( iRow ), ih, w );
-#endif
+#endif //HLTCA_GPUCODE & !HLTCA_GPU_PREFETCHDATA & !HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY
       }
     }
   }
@@ -237,7 +237,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
   &s
 #else
   &/*s*/
-#endif
+#endif //HLTCA_GPUCODE
   , AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam, int iRow )
 {
   // reconstruction of tracklets, tracklets update step
@@ -251,7 +251,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
 
 #ifndef EXTERN_ROW_HITS
   AliHLTTPCCATracklet &tracklet = tracker.Tracklets()[r.fItr];
-#endif
+#endif //EXTERN_ROW_HITS
 
 #ifdef HLTCA_GPU_PREFETCHDATA
   const AliHLTTPCCARow &row = s.fRow[r.fCurrentData];
@@ -259,7 +259,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
   const AliHLTTPCCARow &row = s.fRows[iRow];
 #else
   const AliHLTTPCCARow &row = tracker.Row( iRow );
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
 
   float y0 = row.Grid().YMin();
   float stepY = row.HstepY();
@@ -279,7 +279,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
                  tracklet.SetRowHit(iRow, -1);
 #else
                  tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //EXTERN_ROW_HITS
                  break; // SG!!! - jump over the row
          }
 
@@ -298,7 +298,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
          hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + r.fCurrIH);
 #else
          hh = tracker.HitData(row)[r.fCurrIH];
-#endif
+#endif //HLTCA_GPU_TEXTURE_FETCH
 //#endif
 //#endif
 
@@ -310,7 +310,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
          r.fCurrIH = tex1Dfetch(gAliTexRefs, ((char*) tracker.Data().HitLinkUpData(row) - tracker.Data().GPUTextureBase()) / sizeof(unsigned short) + r.fCurrIH);
 #else
          r.fCurrIH = tracker.HitLinkUpData(row)[r.fCurrIH]; // read from linkup data
-#endif
+#endif //HLTCA_GPU_TEXTURE_FETCH
 //#endif
 
       float x = row.X();
@@ -318,7 +318,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       float z = z0 + hh.y * stepZ;
 #ifdef DRAW
       if ( drawFit ) std::cout << " fit tracklet: new hit " << oldIH << ", xyz=" << x << " " << y << " " << z << std::endl;
-#endif
+#endif //DRAW
 
       if ( iRow == r.fStartRow ) {
         tParam.SetX( x );
@@ -328,7 +328,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         r.fLastZ = z;
         #ifdef DRAW
         if ( drawFit ) std::cout << " fit tracklet " << r.fItr << ", row " << iRow << " first row" << std::endl;
-        #endif
+        #endif //DRAW
       } else {
 
         float err2Y, err2Z;
@@ -354,7 +354,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           std::cout << " fit tracklet " << r.fItr << ", row " << iRow << " transporting.." << std::endl;
           std::cout << " params before transport=" << std::endl;
           tParam.Print();
-          #endif
+          #endif //DRAW
         }
         float sinPhi, cosPhi;
         if ( r.fNHits >= 10 && CAMath::Abs( tParam.SinPhi() ) < .99 ) {
@@ -366,16 +366,16 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         }
         #ifdef DRAW
         if ( drawFit ) std::cout << "sinPhi0 = " << sinPhi << ", cosPhi0 = " << cosPhi << std::endl;
-        #endif
+        #endif //DRAW
         if ( !tParam.TransportToX( x, sinPhi, cosPhi, tracker.Param().ConstBz(), -1 ) ) {
           #ifdef DRAW
           if ( drawFit ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not transport!!" << std::endl;
-                 #endif
+                 #endif //DRAW
 #ifndef EXTERN_ROW_HITS
           tracklet.SetRowHit( iRow, -1 );
 #else
                  tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //EXTERN_ROW_HITS
           break;
         }
         //std::cout<<"mark1 "<<r.fItr<<std::endl;
@@ -390,17 +390,17 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           std::cout << "fit tracklet before filter: " << r.fItr << ", row " << iRow << " errs=" << err2Y << " " << err2Z << std::endl;
           AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 2., 1 );
           AliHLTTPCCADisplay::Instance().Ask();
-                 #endif
+                 #endif //DRAW
         }
         if ( !tParam.Filter( y, z, err2Y, err2Z, .99 ) ) {
           #ifdef DRAW
           if ( drawFit ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not filter!!" << std::endl;
-          #endif
+          #endif //DRAW
 #ifndef EXTERN_ROW_HITS
           tracklet.SetRowHit( iRow, -1 );
 #else
                  tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //EXTERN_ROW_HITS
           break;
         }
       }
@@ -408,14 +408,14 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       tracklet.SetRowHit( iRow, oldIH );
 #else
          tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = oldIH;
-#endif
+#endif //!EXTERN_ROW_HITS
       if ( drawFit ) {
         #ifdef DRAW
         std::cout << "fit tracklet after filter " << r.fItr << ", row " << iRow << std::endl;
         tParam.Print();
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kGreen, 2. );
         AliHLTTPCCADisplay::Instance().Ask();
-               #endif
+               #endif //DRAW
       }
       r.fNHits++;
       r.fLastRow = iRow;
@@ -426,14 +426,14 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
     if ( r.fCurrIH < 0 ) {
       #ifdef DRAW
       if ( drawFitted ) std::cout << "fitted tracklet " << r.fItr << ", nhits=" << r.fNHits << std::endl;
-      #endif
+      #endif //DRAW
       r.fStage = 1;
       //AliHLTTPCCAPerformance::Instance().HNHitsPerSeed()->Fill(r.fNHits);
       if ( r.fNHits < 3 ) { r.fNHits = 0; r.fGo = 0;}//SG!!!
       if ( CAMath::Abs( tParam.SinPhi() ) > .999 ) {
         #ifdef DRAW
         if ( drawFitted ) std::cout << " fitted tracklet  error: sinPhi=" << tParam.SinPhi() << std::endl;
-        #endif
+        #endif //DRAW
         r.fNHits = 0; r.fGo = 0;
       } else {
         //tParam.SetCosPhi( CAMath::Sqrt(1-tParam.SinPhi()*tParam.SinPhi()) );
@@ -444,7 +444,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         tParam.Print();
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue );
         AliHLTTPCCADisplay::Instance().Ask();
-               #endif
+               #endif //DRAW
       }
     }
   } else { // forward/backward searching part
@@ -452,7 +452,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       if ( drawSearch ) {
         #ifdef DRAW
         std::cout << "search tracklet " << r.fItr << " row " << iRow << " miss=" << r.fNMissed << " go=" << r.fGo << " stage=" << r.fStage << std::endl;
-        #endif
+        #endif //DRAW
       }
 
       if ( r.fStage == 2 && ( ( iRow >= r.fEndRow ) ||
@@ -470,17 +470,17 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         #ifdef DRAW
         std::cout << "tracklet " << r.fItr << " before transport to row " << iRow << " : " << std::endl;
         tParam.Print();
-        #endif
+        #endif //DRAW
       }
       if ( !tParam.TransportToX( x, tParam.SinPhi(), tParam.GetCosPhi(), tracker.Param().ConstBz(), .99 ) ) {
         #ifdef DRAW
         if ( drawSearch ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not transport!!" << std::endl;
-        #endif
+        #endif //DRAW
 #ifndef EXTERN_ROW_HITS
                tracklet.SetRowHit(iRow, -1);
 #else
                tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //!EXTERN_ROW_HITS
         break;
       }
       if ( row.NHits() < 1 ) {
@@ -489,7 +489,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
                  tracklet.SetRowHit(iRow, -1);
 #else
                  tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //!EXTERN_ROW_HITS
         break;
       }
       if ( drawSearch ) {
@@ -498,11 +498,11 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         tParam.Print();
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 2., 1 );
         AliHLTTPCCADisplay::Instance().Ask();
-               #endif
+               #endif //DRAW
       }
 #ifdef HLTCA_GPU_PREFETCHDATA
       uint4 *tmpint4 = s.fData[r.fCurrentData];
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
 
 //#ifdef HLTCA_GPU_REORDERHITDATA
 //      const ushort2 *hits = reinterpret_cast<ushort2*>( tmpint4 );
@@ -513,7 +513,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
 //#else
 #ifndef HLTCA_GPU_TEXTURE_FETCH
          const ushort2 *hits = tracker.HitData(row);
-#endif
+#endif //!HLTCA_GPU_TEXTURE_FETCH
 //#endif
 //#endif
 
@@ -537,7 +537,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
 #ifdef DRAW
           std::cout << " tracklet " << r.fItr << ", row " << iRow << ": grid N=" << row.Grid().N() << std::endl;
           std::cout << " tracklet " << r.fItr << ", row " << iRow << ": minbin=" << fIndYmin << std::endl;
-#endif
+#endif //DRAW
         }
         {
           int nY = row.Grid().Ny();
@@ -547,7 +547,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
 //#else
 #ifndef HLTCA_GPU_TEXTURE_FETCH
                  const unsigned short *sGridP = tracker.FirstHitInBin(row);
-#endif
+#endif //!HLTCA_GPU_TEXTURE_FETCH
 //#endif
 
 #ifdef HLTCA_GPU_TEXTURE_FETCH
@@ -560,7 +560,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           fHitYlst = sGridP[fIndYmin+2];
           fHitYfst1 = sGridP[fIndYmin+nY];
           fHitYlst1 = sGridP[fIndYmin+nY+2];
-#endif
+#endif //HLTCA_GPU_TEXTURE_FETCH
           assert( (signed) fHitYfst <= row.NHits() );
           assert( (signed) fHitYlst <= row.NHits() );
           assert( (signed) fHitYfst1 <= row.NHits() );
@@ -580,21 +580,21 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
               }
               std::cout << std::endl;
             }
-#endif
+#endif //DRAW
           }
 #ifdef DRAW
           if ( sGridP[row.Grid().N()] != row.NHits() ) {
             std::cout << " grid, row " << iRow << ": nHits=" << row.NHits() << ", grid n=" << row.Grid().N() << ", c[n]=" << sGridP[row.Grid().N()] << std::endl;
             //exit(0);
           }
-#endif
+#endif //DRAW
         }
 #ifdef DRAW
         if ( drawSearch ) {
           std::cout << " tracklet " << r.fItr << ", row " << iRow << ", yz= " << fY << "," << fZ << ": search hits=" << fHitYfst << " " << fHitYlst << " / " << fHitYfst1 << " " << fHitYlst1 << std::endl;
           std::cout << " hit search :" << std::endl;
         }
-#endif
+#endif //DRAW
         for ( unsigned int fIh = fHitYfst; fIh < fHitYlst; fIh++ ) {
           assert( (signed) fIh < row.NHits() );
           ushort2 hh;
@@ -602,14 +602,14 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
                 hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
 #else
                  hh = hits[fIh];
-#endif
+#endif //HLTCA_GPU_TEXTURE_FETCH
           int ddy = ( int )( hh.x ) - fY0;
           int ddz = ( int )( hh.y ) - fZ0;
           int dds = CAMath::Abs( ddy ) + CAMath::Abs( ddz );
           if ( drawSearch ) {
             #ifdef DRAW
             std::cout << fIh << ": hityz= " << hh.x << " " << hh.y << "(" << hh.x*stepY << " " << hh.y*stepZ << "), trackyz=" << fY0 << " " << fZ0 << "(" << fY0*stepY << " " << fZ0*stepZ << "), dy,dz,ds= " << ddy << " " << ddz << " " << dds << "(" << ddy*stepY << " " << ddz*stepZ << std::endl;
-            #endif
+            #endif //DRAW
           }
           if ( dds < ds ) {
             ds = dds;
@@ -623,14 +623,14 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
                  hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
 #else
                  hh = hits[fIh];
-#endif
+#endif //HLTCA_GPU_TEXTURE_FETCH
           int ddy = ( int )( hh.x ) - fY0;
           int ddz = ( int )( hh.y ) - fZ0;
           int dds = CAMath::Abs( ddy ) + CAMath::Abs( ddz );
           if ( drawSearch ) {
             #ifdef DRAW
             std::cout << fIh << ": hityz= " << hh.x << " " << hh.y << "(" << hh.x*stepY << " " << hh.y*stepZ << "), trackyz=" << fY0 << " " << fZ0 << "(" << fY0*stepY << " " << fZ0*stepZ << "), dy,dz,ds= " << ddy << " " << ddz << " " << dds << "(" << ddy*stepY << " " << ddz*stepZ << std::endl;
-            #endif
+            #endif //DRAW
           }
           if ( dds < ds ) {
             ds = dds;
@@ -645,7 +645,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
                  tracklet.SetRowHit(iRow, -1);
 #else
                  tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //!EXTERN_ROW_HITS
                  break;
          }
       if ( drawSearch ) {
@@ -655,7 +655,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         AliHLTTPCCADisplay::Instance().Ask();
         AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best, kWhite, 1 );
         AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best );
-               #endif
+               #endif //DRAW
       }
 
       ushort2 hh;
@@ -663,7 +663,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
                 hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + best);
 #else
                  hh = hits[best];
-#endif
+#endif //HLTCA_GPU_TEXTURE_FETCH
 
       //std::cout<<"mark 3, "<<r.fItr<<std::endl;
       //tParam.Print();
@@ -685,19 +685,19 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         #ifdef DRAW
         std::cout << "dy,sy= " << dy << " " << CAMath::Sqrt( sy2 ) << ", dz,sz= " << dz << " " << CAMath::Sqrt( sz2 ) << std::endl;
         std::cout << "dy,dz= " << dy << " " << dz << ", sy,sz= " << CAMath::Sqrt( sy2 ) << " " << CAMath::Sqrt( sz2 ) << ", sy,sz= " << CAMath::Sqrt( kFactor*( tParam.GetErr2Y() +  err2Y ) ) << " " << CAMath::Sqrt( kFactor*( tParam.GetErr2Z() +  err2Z ) ) << std::endl;
-        #endif
+        #endif //DRAW
       }
       if ( CAMath::FMulRZ( dy, dy ) > sy2 || CAMath::FMulRZ( dz, dz ) > sz2  ) {
         if ( drawSearch ) {
           #ifdef DRAW
           std::cout << "found hit is out of the chi2 window\n " << std::endl;
-          #endif
+          #endif //DRAW
         }
 #ifndef EXTERN_ROW_HITS
                tracklet.SetRowHit(iRow, -1);
 #else
                tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
-#endif
+#endif //!EXTERN_ROW_HITS
         break;
       }
 #ifdef DRAW
@@ -705,12 +705,12 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       //std::cout<<"hit search before filter: "<<r.fItr<<", row "<<iRow<<std::endl;
       //AliHLTTPCCADisplay::Instance().DrawTracklet(tParam, hitstore, kBlue);
       //AliHLTTPCCADisplay::Instance().Ask();
-#endif
+#endif //DRAW
       if ( !tParam.Filter( y, z, err2Y, err2Z, .99 ) ) {
         if ( drawSearch ) {
           #ifdef DRAW
           std::cout << "tracklet " << r.fItr << " at row " << iRow << " : can not filter!!!! " << std::endl;
-          #endif
+          #endif //DRAW
         }
         break;
       }
@@ -718,14 +718,14 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
          tracklet.SetRowHit( iRow, best );
 #else
          tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = best;
-#endif
+#endif //!EXTERN_ROW_HITS
       if ( drawSearch ) {
         #ifdef DRAW
         std::cout << "tracklet " << r.fItr << " after filter at row " << iRow << " : " << std::endl;
         tParam.Print();
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kRed );
         AliHLTTPCCADisplay::Instance().Ask();
-               #endif
+               #endif //DRAW
       }
       r.fNHits++;
       r.fNMissed = 0;
@@ -809,7 +809,7 @@ GPUd() int AliHLTTPCCATrackletConstructor::FetchTracklet(AliHLTTPCCATracker &tra
                                        sMem.fNextTrackletNoDummy = 1;
                                }
                        }
-#endif
+#endif //HLTCA_GPU_SCHED_FIXED_START
                }
                else
                {
@@ -845,7 +845,7 @@ GPUd() int AliHLTTPCCATrackletConstructor::FetchTracklet(AliHLTTPCCATracker &tra
        {
                return(-1);
        }
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_NMEMTHREDS > 0
        else if (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS >= sMem.fNextTrackletCount)
        {
                return(-1);             //No track in this RowBlock for this thread
@@ -889,7 +889,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
        {
                pTracker[0].BlockStartingTracklet()[i].x = pTracker[0].BlockStartingTracklet()[i].y = 0;
        }
-#endif
+#endif //HLTCA_GPU_EMULATION_SINGLE_TRACKLET
 
        GPUshared() AliHLTTPCCASharedMemory sMem;
 
@@ -899,7 +899,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                sMem.fNextTrackletFirstRun = 1;
        }
        __syncthreads();
-#endif
+#endif //HLTCA_GPU_SCHED_FIXED_START
 
 #ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
        if (threadIdx.x == TRACKLET_CONSTRUCTOR_NMEMTHREDS)
@@ -907,7 +907,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                sMem.fMaxSync = 0;
        }
        int threadSync = 0;
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
 
        for (int iReverse = 0;iReverse < 2;iReverse++)
        {
@@ -917,7 +917,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                        int iSlice = tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + (HLTCA_GPU_BLOCK_COUNT % tracker.GPUParametersConst()->fGPUnSlices != 0 && tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + 1) % HLTCA_GPU_BLOCK_COUNT != 0)) / HLTCA_GPU_BLOCK_COUNT;
 #else
                        for (int iSlice = 0;iSlice < pTracker[0].GPUParametersConst()->fGPUnSlices;iSlice++)
-#endif
+#endif //HLTCA_GPU_SCHED_FIXED_SLICE
                        {
                                AliHLTTPCCATracker &tracker = pTracker[iSlice];
                                if (sMem.fNextTrackletFirstRun && iSlice != tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + (HLTCA_GPU_BLOCK_COUNT % tracker.GPUParametersConst()->fGPUnSlices != 0 && tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + 1) % HLTCA_GPU_BLOCK_COUNT != 0)) / HLTCA_GPU_BLOCK_COUNT)
@@ -938,7 +938,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                        CAMath::AtomicMax(&sMem.fMaxSync, threadSync);
                                        __syncthreads();
                                        threadSync = CAMath::Min(sMem.fMaxSync, 100000000 / blockDim.x / gridDim.x);
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
 #ifndef HLTCA_GPU_PREFETCHDATA
                                        if (!sharedRowsInitialized)
                                        {
@@ -962,10 +962,10 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                                {
                                                        reinterpret_cast<int*>(&sMem.fRows)[i] = reinterpret_cast<int*>(tracker.SliceDataRows())[i];
                                                }
-#endif
+#endif //HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY
                                                sharedRowsInitialized = 1;
                                        }
-#endif
+#endif //!HLTCA_GPU_PREFETCHDATA
 #ifdef HLTCA_GPU_RESCHED
                                        short2 storeToRowBlock;
                                        int storePosition = 0;
@@ -975,7 +975,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                                const int nRowBlock = (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS) % (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1);
                                                sMem.fTrackletStoreCount[nReverse][nRowBlock] = 0;
                                        }
-#endif
+#endif //HLTCA_GPU_RESCHED
                                        __syncthreads();
                                        AliHLTTPCCATrackParam tParam;
                                        AliHLTTPCCAThreadMemory rMem;
@@ -987,7 +987,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                        {
                                                tracker.GPUParameters()->fGPUError = 1;
                                        }
-#endif
+#endif //HLTCA_GPU_EMULATION_DEBUG_TRACKLET
                                        AliHLTTPCCAThreadMemory &rMemGlobal = tracker.GPUTrackletTemp()[iTracklet].fThreadMem;
                                        AliHLTTPCCATrackParam &tParamGlobal = tracker.GPUTrackletTemp()[iTracklet].fParam;
                                        if (mustInit)
@@ -1011,7 +1011,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                        {
                                                ReadData(threadIdx.x, sMem, rMem, tracker, iReverse ? (HLTCA_ROW_COUNT - 1 - iRowBlock * HLTCA_GPU_SCHED_ROW_STEP) : (CAMath::Max(1, iRowBlock * HLTCA_GPU_SCHED_ROW_STEP)));
                                        }
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
                                        rMem.fItr = iTracklet;
                                        rMem.fGo = (iTracklet >= 0);
 
@@ -1021,7 +1021,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
 #ifdef HLTCA_GPU_PREFETCHDATA
                                        rMem.fCurrentData ^= 1;
                                        __syncthreads();
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
                                        if (iReverse)
                                        {
                                                for (int j = HLTCA_ROW_COUNT - 1 - iRowBlock * HLTCA_GPU_SCHED_ROW_STEP;j >= CAMath::Max(0, HLTCA_ROW_COUNT - (iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP);j--)
@@ -1029,14 +1029,14 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
 #ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
                                                        if (rMem.fNMissed <= kMaxRowGap && rMem.fGo && !(j >= rMem.fEndRow || ( j >= rMem.fStartRow && j - rMem.fStartRow % 2 == 0)))
                                                                pTracker[0].fStageAtSync[threadSync++ * blockDim.x * gridDim.x + blockIdx.x * blockDim.x + threadIdx.x] = rMem.fStage + 1;
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
 #ifdef HLTCA_GPU_PREFETCHDATA
                                                        if (threadIdx.x < TRACKLET_CONSTRUCTOR_NMEMTHREDS && j > CAMath::Max(0, HLTCA_ROW_COUNT - (iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP))
                                                        {
                                                                ReadData(threadIdx.x, sMem, rMem, tracker, j - 1);
                                                        }
                                                        else
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
                                                        if (iTracklet >= 0)
                                                        {
                                                                UpdateTracklet(gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);
@@ -1045,13 +1045,13 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                                                        rMem.fGo = 0;
 #ifndef HLTCA_GPU_PREFETCHDATA
                                                                        break;
-#endif
+#endif //!HLTCA_GPU_PREFETCHDATA
                                                                }
                                                        }
 #ifdef HLTCA_GPU_PREFETCHDATA
                                                        __syncthreads();
                                                        rMem.fCurrentData ^= 1;
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
                                                }
                                                        
                                                if (iTracklet >= 0 && (!rMem.fGo || iRowBlock == HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP))
@@ -1066,25 +1066,25 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
 #ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
                                                        if (rMem.fNMissed <= kMaxRowGap && rMem.fGo && j >= rMem.fStartRow && (rMem.fStage > 0 || rMem.fCurrIH >= 0 || (j - rMem.fStartRow) % 2 == 0 ))
                                                                pTracker[0].fStageAtSync[threadSync++ * blockDim.x * gridDim.x + blockIdx.x * blockDim.x + threadIdx.x] = rMem.fStage + 1;
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
 #ifdef HLTCA_GPU_PREFETCHDATA
                                                        if (threadIdx.x < TRACKLET_CONSTRUCTOR_NMEMTHREDS && j < CAMath::Min((iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP, HLTCA_ROW_COUNT) - 1)
                                                        {
                                                                ReadData(threadIdx.x, sMem, rMem, tracker, j + 1);
                                                        }
                                                        else
-#endif 
+#endif //HLTCA_GPU_PREFETCHDATA
                                                        if (iTracklet >= 0)
                                                        {
                                                                UpdateTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);
 #ifndef HLTCA_GPU_PREFETCHDATA
                                                                //if (rMem.fNMissed > kMaxRowGap || rMem.fGo == 0) break;       //DR!!! CUDA Crashes with this enabled
-#endif
+#endif //!HLTCA_GPU_PREFETCHDATA
                                                        }
 #ifdef HLTCA_GPU_PREFETCHDATA
                                                        __syncthreads();
                                                        rMem.fCurrentData ^= 1;
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
                                                }
                                                if (rMem.fGo && (rMem.fNMissed > kMaxRowGap || iRowBlock == HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP))
                                                {
@@ -1092,7 +1092,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                                        if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) )
 #else
                                                        if ( !tParam.TransportToX( sMem.fRows[ rMem.fEndRow ].X(), tracker.Param().ConstBz(), .999 ) )
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA | !HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY
                                                        {
                                                                rMem.fGo = 0;
                                                        }
@@ -1140,7 +1140,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
 
                                                StoreTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam );
                                        }
-#endif
+#endif //HLTCA_GPU_RESCHED
 #ifdef HLTCA_GPU_RESCHED
                                        __syncthreads();
                                        if (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS < 2 * (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1))
@@ -1158,14 +1158,11 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
                                                tracker.RowBlockTracklets(storeToRowBlock.y, storeToRowBlock.x)[sMem.fTrackletStoreCount[storeToRowBlock.y][storeToRowBlock.x] + storePosition] = iTracklet;
                                        }
                                        __syncthreads();
-#endif
+#endif //HLTCA_GPU_RESCHED
                                }
                        }
                }
        }
-#ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
-
-#endif
 }
 
 GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorInit(int iTracklet, AliHLTTPCCATracker &tracker)
@@ -1177,7 +1174,7 @@ AliHLTTPCCAHitId id = tracker.TrackletStartHits()[iTracklet];
 #ifdef HLTCA_GPU_SCHED_FIXED_START
        const int firstDynamicTracklet = tracker.GPUParameters()->fScheduleFirstDynamicTracklet;
        if (iTracklet >= firstDynamicTracklet)
-#endif
+#endif //HLTCA_GPU_SCHED_FIXED_START
        {
                const int firstTrackletInRowBlock = CAMath::Max(firstDynamicTracklet, tracker.RowStartHitCountOffset()[CAMath::Max(id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP * HLTCA_GPU_SCHED_ROW_STEP, 1)].z);
                if (iTracklet == firstTrackletInRowBlock)
@@ -1194,7 +1191,7 @@ AliHLTTPCCAHitId id = tracker.TrackletStartHits()[iTracklet];
                }
                tracker.RowBlockTracklets(0, id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP)[iTracklet - firstTrackletInRowBlock] = iTracklet;
        }
-#endif
+#endif //!HLTCA_GPU_EMULATION_SINGLE_TRACKLET
 }
 
 GPUg() void AliHLTTPCCATrackletConstructorInit(int iSlice)
@@ -1213,7 +1210,8 @@ GPUg() void AliHLTTPCCATrackletConstructorNewGPU()
        AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU(pTracker);
 }
 
-#else
+#else //HLTCA_GPUCODE
+
 GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewCPU(AliHLTTPCCATracker &tracker)
 {
        //Tracklet constructor simple CPU Function that does not neew a scheduler
@@ -1259,4 +1257,4 @@ GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewCPU
                StoreTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam );
        }
 }
-#endif
+#endif //HLTCA_GPUCODE
index dfd272c..b8b40c3 100644 (file)
@@ -13,7 +13,7 @@
 #define HLTCA_GPU_USE_INT short
 #else
 #define HLTCA_GPU_USE_INT int
-#endif
+#endif //HLTCA_GPUCODE
 
 #include "AliHLTTPCCADef.h"
 #include "AliHLTTPCCAGPUConfig.h"
@@ -40,7 +40,7 @@ class AliHLTTPCCATrackletConstructor
         AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/ )
                        : fNextTrackletFirst(0), fNextTrackletCount(0), fNextTrackletNoDummy(0), fNextTrackletStupidDummy(0), fNextTrackletFirstRun(0), fNTracklets(0), fSliceDone(0) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
-#endif
+#endif //HLTCA_GPUCODE
 
       protected:
 #ifdef HLTCA_GPU_PREFETCHDATA
@@ -48,7 +48,7 @@ class AliHLTTPCCATrackletConstructor
                AliHLTTPCCARow fRow[2]; // row
 #else
                AliHLTTPCCARow fRows[HLTCA_ROW_COUNT];
-#endif
+#endif //HLTCA_GPU_PREFETCHDATA
                int fNextTrackletFirst;
                int fNextTrackletCount;
                int fNextTrackletNoDummy;
@@ -59,7 +59,7 @@ class AliHLTTPCCATrackletConstructor
 
 #ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
                int fMaxSync;
-#endif
+#endif //HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
 
                int fTrackletStoreCount[2][HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1];
     };
@@ -75,7 +75,7 @@ class AliHLTTPCCATrackletConstructor
         AliHLTTPCCAThreadMemory( const AliHLTTPCCAThreadMemory& /*dummy*/ )
             : fItr( 0 ), fFirstRow( 0 ), fLastRow( 0 ), fStartRow( 0 ), fEndRow( 0 ), fCurrIH( 0 ), fGo( 0 ), fCurrentData( 0 ), fStage( 0 ), fNHits( 0 ), fNMissed( 0 ), fLastY( 0 ), fLastZ( 0 ) {}
         AliHLTTPCCAThreadMemory& operator=( const AliHLTTPCCAThreadMemory& /*dummy*/ ) { return *this; }
-#endif
+#endif //!HLTCA_GPUCODE
 
       protected:
         int fItr; // track index
@@ -118,7 +118,7 @@ class AliHLTTPCCATrackletConstructor
        GPUd() static void CopyTrackletTempData( AliHLTTPCCAThreadMemory &rMemSrc, AliHLTTPCCAThreadMemory &rMemDst, AliHLTTPCCATrackParam &tParamSrc, AliHLTTPCCATrackParam &tParamDst);
 #else
        GPUd() static void AliHLTTPCCATrackletConstructorNewCPU(AliHLTTPCCATracker &tracker);
-#endif
+#endif //HLTCA_GPUCODE
 
     GPUd() static bool SAVE() { return 1; }
 
@@ -128,10 +128,8 @@ class AliHLTTPCCATrackletConstructor
 #else
     //GPUhd() inline int NMemThreads() { return 1; }
 #define TRACKLET_CONSTRUCTOR_NMEMTHREDS 1
-#endif
+#endif //!HLTCA_GPUCODE
 
 };
 
-
-
-#endif
+#endif //ALIHLTTPCCATRACKLETCONSTRUCTOR_H
index 70dde8c..7e689ea 100644 (file)
@@ -48,7 +48,7 @@ GPUd() void AliHLTTPCCATrackletSelector::Thread
        {
                tracker.GPUParameters()->fGPUError = 1;
        }
-#endif
+#endif //HLTCA_GPU_EMULATION_DEBUG_TRACKLET
 
          while (tracker.Tracklets()[itr].NHits() == 0)
          {
@@ -90,7 +90,7 @@ GPUd() void AliHLTTPCCATrackletSelector::Thread
         int ih = tracker.TrackletRowHits()[irow * s.fNTracklets + itr];
 #else
                int ih = tracklet.RowHit( irow );
-#endif
+#endif //EXTERN_ROW_HITS
         if ( ih >= 0 ) {
           const AliHLTTPCCARow &row = tracker.Row( irow );
           bool own = ( tracker.HitWeight( row, ih ) <= w );
@@ -101,7 +101,7 @@ GPUd() void AliHLTTPCCATrackletSelector::Thread
                        if (nHits < HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE)
                                s.fHits[iThread][nHits].Set( irow, ih );
                        else
-#endif
+#endif //HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE != 0
                                trackHits[nHits - HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE].Set( irow, ih );
             nHits++;
             if ( !own ) nShared++;
@@ -118,7 +118,7 @@ GPUd() void AliHLTTPCCATrackletSelector::Thread
                                CAMath::AtomicExch( tracker.NTracks(), 0 );
                                return;
                        }
-#endif
+#endif //HLTCA_GPUCODE
             nFirstTrackHit = CAMath::AtomicAdd( tracker.NTrackHits(), nHits );
                        tracker.Tracks()[itrout].SetAlive(1);
                        tracker.Tracks()[itrout].SetParam(tracklet.Param());
@@ -129,7 +129,7 @@ GPUd() void AliHLTTPCCATrackletSelector::Thread
                                if (jh < HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE)
                                        tracker.TrackHits()[nFirstTrackHit + jh] = s.fHits[iThread][jh];
                                else
-#endif
+#endif //HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE != 0
                                        tracker.TrackHits()[nFirstTrackHit + jh] = trackHits[jh - HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE];
             }
           }
index a08f452..0b7e527 100644 (file)
@@ -32,7 +32,7 @@ class AliHLTTPCCATrackletSelector
         int fNTracklets; // n of tracklets
 #if HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE != 0
                AliHLTTPCCAHitId fHits[HLTCA_GPU_THREAD_COUNT][HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE];
-#endif
+#endif //HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE != 0
        };
 
     GPUd() static int NThreadSyncPoints() { return 1; }
@@ -43,4 +43,4 @@ class AliHLTTPCCATrackletSelector
 };
 
 
-#endif
+#endif //ALIHLTTPCCATRACKLETSELECTOR_H
index d9d4472..32822b8 100644 (file)
@@ -62,6 +62,6 @@ class AliTPCtrackerCA : public AliTracker
 };
 
 
-#endif
+#endif //ALITPCTRACKERCA_H
 
 
index 2a48f76..5fe6692 100644 (file)
@@ -19,7 +19,7 @@
 
 #ifndef assert
 #include <assert.h>
-#endif
+#endif //!assert
 
 template<unsigned int X>
 GPUhd() static inline void AlignTo( char *&mem )
diff --git a/HLT/libAliHLTTPC.conf b/HLT/libAliHLTTPC.conf
deleted file mode 100644 (file)
index 04f6ec2..0000000
+++ /dev/null
@@ -1,12 +0,0 @@
-#-*- Mode: Makefile -*-
-# $Id$
-#GPU Tracker Build for the libAliHLTTPC
-
-ifdef NVCC
-ELIBS          += cudart cuda
-ELIBSDIR       += $(NVCC:/bin/nvcc=/lib64)
-MODULE_CUHDRS  += $(TRACKING_CA) tracking-ca/AliHLTTPCCAGPUTracker.h
-MODULE_CUSRCS  += tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cu
-EDEFINE                += -DBUILD_GPU
-endif
-