updae of the GPU framework
authorsgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Sat, 31 Oct 2009 07:29:40 +0000 (07:29 +0000)
committersgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Sat, 31 Oct 2009 07:29:40 +0000 (07:29 +0000)
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cu
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceData.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceData.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerFramework.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx

index a9c8f45ac71f96fe0bb098d9cea028c38dba2a88..cb2a3f8869fb8b00363fca072fd5f04dd0b36a0d 100644 (file)
@@ -28,9 +28,9 @@ int AliHLTTPCCAGPUTracker::InitGPU(int /*sliceCount*/, int /*forceDeviceID*/)
     return(1);
 }
 void AliHLTTPCCAGPUTracker::StandalonePerfTime(int /*iSlice*/, int /*i*/) {}
-//template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int alignment) {return(NULL);}
+template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int alignment) {return(NULL);}
 //bool AliHLTTPCCAGPUTracker::CudaFailedMsg(cudaError_t error) {return(true);}
-//int AliHLTTPCCAGPUTracker::CUDASync() {return(1);}
+int AliHLTTPCCAGPUTracker::CUDASync(char* /*text*/) {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);}
@@ -38,4 +38,6 @@ 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
index 3c213968fb5536455f72d0aeb8436c843b504565..6d5f4ed2c670d406ef4da91e0d0b482027e10782 100644 (file)
@@ -59,6 +59,8 @@ private:
 
        void DumpRowBlocks(AliHLTTPCCATracker* tracker, int iSlice, bool check = true);
        int GetThread();
+       void ReleaseGlobalLock(void* sem);
+       int CheckMemorySizes(int sliceCount);
 
        AliHLTTPCCATracker *fGpuTracker;
        void* fGPUMemory;
index 9faf18c74008b85876cc8a39ec35be4766ac1435..9d272d53194a48b9640107c49e92ce83034497de 100644 (file)
 
 #include <cuda.h>
 #ifdef R__WIN32
-
 #else
 #include <sys/syscall.h>
+#include <semaphore.h>
+#include <fcntl.h>
 #endif
 
 #include "AliHLTTPCCADef.h"
@@ -80,13 +81,75 @@ ClassImp( AliHLTTPCCAGPUTracker )
 
 bool AliHLTTPCCAGPUTracker::fgGPUUsed = false;
 
+#define SemLockName "AliceHLTTPCCAGPUTrackerInitLockSem"
+
+void AliHLTTPCCAGPUTracker::ReleaseGlobalLock(void* sem)
+{
+       //Release the global named semaphore that locks GPU Initialization
+#ifdef R__WIN32
+       HANDLE* h = (HANDLE*) sem;
+       ReleaseSemaphore(*h, 1, NULL);
+       CloseHandle(*h);
+       delete h;
+#else
+       sem_t* pSem = (sem_t*) sem;
+       sem_post(pSem);
+       sem_unlink(SemLockName);
+#endif
+}
+
+int AliHLTTPCCAGPUTracker::CheckMemorySizes(int sliceCount)
+{
+       //Check constants for correct memory sizes
+  if (sizeof(AliHLTTPCCATracker) * sliceCount > HLTCA_GPU_TRACKER_OBJECT_MEMORY)
+  {
+         HLTError("Insufficiant Tracker Object Memory");
+         return(1);
+  }
+
+  if (fgkNSlices * AliHLTTPCCATracker::CommonMemorySize() > HLTCA_GPU_COMMON_MEMORY)
+  {
+         HLTError("Insufficiant Common Memory");
+         return(1);
+  }
+
+  if (fgkNSlices * (HLTCA_ROW_COUNT + 1) * sizeof(AliHLTTPCCARow) > HLTCA_GPU_ROWS_MEMORY)
+  {
+         HLTError("Insufficiant Row Memory");
+         return(1);
+  }
+  return(0);
+}
+
 int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
 {
        //Find best CUDA device, initialize and allocate memory
-       
+
+       if (CheckMemorySizes(sliceCount)) return(1);
+
+#ifdef R__WIN32
+       HANDLE* semLock = new HANDLE;
+       *semLock = CreateSemaphore(NULL, 1, 1, SemLockName);
+       if (*semLock == NULL)
+       {
+               HLTError("Error creating GPUInit Semaphore");
+               return(1);
+       }
+       WaitForSingleObject(*semLock, INFINITE);
+#else
+       sem_t* semLock = sem_open(SemLockName, O_CREAT, 0x01B6, 1);
+       if (semLock == SEM_FAILED)
+       {
+               HLTError("Error creating GPUInit Semaphore");
+               return(1);
+       }
+       sem_wait(semLock);
+#endif
+
        if (fgGPUUsed)
        {
            HLTWarning("CUDA already used by another AliHLTTPCCAGPUTracker running in same process");
+               ReleaseGlobalLock(semLock);
            return(1);
        }
        fgGPUUsed = 1;
@@ -103,6 +166,7 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
        {
                HLTError("Error getting CUDA Device Count");
                fgGPUUsed = 0;
+               ReleaseGlobalLock(semLock);
                return(1);
        }
        if (fDebugLevel >= 2) HLTInfo("Available CUDA devices:");
@@ -132,6 +196,7 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
        {
                HLTWarning("No CUDA Device available, aborting CUDA Initialisation");
                fgGPUUsed = 0;
+               ReleaseGlobalLock(semLock);
                return(1);
        }
 
@@ -168,36 +233,15 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
   {
        HLTError( "Unsupported CUDA Device" );
        fgGPUUsed = 0;
+       ReleaseGlobalLock(semLock);
        return(1);
   }
 
-  if (sizeof(AliHLTTPCCATracker) * sliceCount > HLTCA_GPU_TRACKER_OBJECT_MEMORY)
-  {
-         HLTError("Insufficiant Tracker Object Memory");
-         fgGPUUsed = 0;
-         return(1);
-  }
-  
   if (CudaFailedMsg(cudaSetDevice(cudaDevice)))
   {
          HLTError("Could not set CUDA Device!");
          fgGPUUsed = 0;
-         return(1);
-  }
-
-  if (fgkNSlices * AliHLTTPCCATracker::CommonMemorySize() > HLTCA_GPU_COMMON_MEMORY)
-  {
-         HLTError("Insufficiant Common Memory");
-         cudaThreadExit();
-         fgGPUUsed = 0;
-         return(1);
-  }
-
-  if (fgkNSlices * (HLTCA_ROW_COUNT + 1) * sizeof(AliHLTTPCCARow) > HLTCA_GPU_ROWS_MEMORY)
-  {
-         HLTError("Insufficiant Row Memory");
-         cudaThreadExit();
-         fgGPUUsed = 0;
+         ReleaseGlobalLock(semLock);
          return(1);
   }
 
@@ -206,8 +250,10 @@ int AliHLTTPCCAGPUTracker::InitGPU(int sliceCount, int forceDeviceID)
          HLTError("CUDA Memory Allocation Error");
          cudaThreadExit();
          fgGPUUsed = 0;
+         ReleaseGlobalLock(semLock);
          return(1);
   }
+  ReleaseGlobalLock(semLock);
   if (fDebugLevel >= 1) HLTInfo("GPU Memory used: %d", (int) fGPUMemSize);
   int hostMemSize = HLTCA_GPU_ROWS_MEMORY + HLTCA_GPU_COMMON_MEMORY + sliceCount * (HLTCA_GPU_SLICE_DATA_MEMORY + HLTCA_GPU_TRACKS_MEMORY) + HLTCA_GPU_TRACKER_OBJECT_MEMORY;
   if (CudaFailedMsg(cudaMallocHost(&fHostLockedMemory, hostMemSize)))
@@ -619,10 +665,8 @@ int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTT
                if (fDebugLevel >= 4)
                {
                        if (fDebugLevel >= 5) HLTInfo("Allocating Debug Output Memory");
-                       fSlaveTrackers[firstSlice + iSlice].TrackletMemory() = reinterpret_cast<char*> ( new uint4 [ fGpuTracker[iSlice].TrackletMemorySize()/sizeof( uint4 ) + 100] );
-                       fSlaveTrackers[firstSlice + iSlice].SetPointersTracklets( HLTCA_GPU_MAX_TRACKLETS );
-                       fSlaveTrackers[firstSlice + iSlice].HitMemory() = reinterpret_cast<char*> ( new uint4 [ fGpuTracker[iSlice].HitMemorySize()/sizeof( uint4 ) + 100] );
-                       fSlaveTrackers[firstSlice + iSlice].SetPointersHits( pClusterData[iSlice].NumberOfClusters() );
+                       fSlaveTrackers[firstSlice + iSlice].SetGPUTrackerTrackletsMemory(reinterpret_cast<char*> ( new uint4 [ fGpuTracker[iSlice].TrackletMemorySize()/sizeof( uint4 ) + 100] ), HLTCA_GPU_MAX_TRACKLETS);
+                       fSlaveTrackers[firstSlice + iSlice].SetGPUTrackerHitsMemory(reinterpret_cast<char*> ( new uint4 [ fGpuTracker[iSlice].HitMemorySize()/sizeof( uint4 ) + 100]), pClusterData[iSlice].NumberOfClusters() );
                }
                
                if (CUDASync("Initialization")) return(1);
@@ -939,6 +983,7 @@ int AliHLTTPCCAGPUTracker::ExitGPU()
 
 void AliHLTTPCCAGPUTracker::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val)
 {
+       //Set Output Control Pointers
        fOutputControl = val;
        for (int i = 0;i < fgkNSlices;i++)
        {
@@ -948,6 +993,7 @@ void AliHLTTPCCAGPUTracker::SetOutputControl( AliHLTTPCCASliceOutput::outputCont
 
 int AliHLTTPCCAGPUTracker::GetThread()
 {
+       //Get Thread ID
 #ifdef R__WIN32
        return((int) (size_t) GetCurrentThread());
 #else
index b9cb28410aa5dc1685cb2441e1f7e99893b5eb93..a1bf697b1174e54480be471e8b19972e51675e4a 100644 (file)
@@ -174,7 +174,6 @@ GPUd() void AliHLTTPCCANeighboursFinder::Thread
 #endif
 
         AliHLTTPCCAHitArea areaDn, areaUp;
-        // TODO: for NVIDIA GPUs it should use the GridContentUp/-Dn that got copied into shared mem
         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 );
 
index 3175bae880b9784ec1878e5f0cd38fc39ee6f1fb..7100ad8dd5c3415048f8ac3825526786d249fd72 100644 (file)
@@ -67,7 +67,7 @@ inline void AliHLTTPCCASliceData::CreateGrid( AliHLTTPCCARow *row, const AliHLTT
                      CAMath::Max( ( zMax - zMin ) * norm, 2.f ) );
 }
 
-inline void AliHLTTPCCASliceData::PackHitData( AliHLTTPCCARow *row, const AliHLTArray<AliHLTTPCCAHit> &binSortedHits )
+inline void AliHLTTPCCASliceData::PackHitData( AliHLTTPCCARow* const row, const AliHLTArray<AliHLTTPCCAHit> &binSortedHits )
 {
   // hit data packing
 
index c6fee6e193b4a68024fb5fcb1ef4168fc8e43fb5..59e4bbc0c9eea71334df66cd916267cbccfb1ad4 100644 (file)
@@ -148,7 +148,7 @@ class AliHLTTPCCASliceData
        GPUh() size_t GpuMemorySize() const {return(fGpuMemorySize); }
        GPUh() int* HitWeights() const {return(fHitWeights); }
 
-       GPUhd() void SetGPUTextureBase(char* val) {fGPUTextureBase = val;}
+       GPUhd() void SetGPUTextureBase(char* const val) {fGPUTextureBase = val;}
        GPUhd() char* GPUTextureBase() const { return(fGPUTextureBase); }
        GPUhd() char* GPUTextureBaseConst() const { return(fGPUTextureBase); }
 
index e5a0f66427c6b4e7fee5c8ebfbbf4446ef20c1cf..59a1bda3b49890988c0b1af70c3b52869a9f7203 100644 (file)
@@ -77,7 +77,7 @@ class AliHLTTPCCATrackParam
     GPUd() const float *GetPar() const { return fParam.GetPar(); }
        GPUd() float GetPar(int i) const { return(fParam.GetPar(i)); }
     GPUd() const float *GetCov() const { return fC; }
-       GPUd() float GetCov(int i) {return fC[i]; }
+       GPUd() float GetCov(int i) const {return fC[i]; }
 
     GPUhd() void SetPar( int i, float v ) { fParam.SetPar(i, v); }
     GPUhd() void SetCov( int i, float v ) { fC[i] = v; }
@@ -137,6 +137,9 @@ class AliHLTTPCCATrackParam
 
     GPUd() void Print() const;
 
+#ifndef HLTCA_GPUCODE
+  private:
+#endif
        AliHLTTPCCATrackParam2 fParam; // Track Parameters
 
   private:
index fa9866ca14d9509a96b334e26a24c1f908f0e4b3..cf675f12120eec6ae0aafcf88250d4d1847d1a88 100644 (file)
@@ -220,15 +220,14 @@ class AliHLTTPCCATracker
     return fData.HitWeight( row, hitIndex );
   }
   
-  GPUhd() int NTracklets() const { return fCommonMem->fNTracklets; }
-  GPUhd() int  *NTracklets() { return &fCommonMem->fNTracklets; }
+  GPUhd() int *NTracklets() const { return &fCommonMem->fNTracklets; }
   
   GPUhd() const AliHLTTPCCAHitId &TrackletStartHit( int i ) const { return fTrackletStartHits[i]; }
   GPUhd() AliHLTTPCCAHitId *TrackletStartHits() const { return fTrackletStartHits; }
   GPUhd() AliHLTTPCCAHitId *TrackletTmpStartHits() const { return fTrackletTmpStartHits; }
   GPUhd() const AliHLTTPCCATracklet &Tracklet( int i ) const { return fTracklets[i]; }
   GPUhd() AliHLTTPCCATracklet  *Tracklets() const { return fTracklets;}
-  GPUhd() int* TrackletRowHits() { return fTrackletRowHits; }
+  GPUhd() int* TrackletRowHits() const { return fTrackletRowHits; }
 
   GPUhd() int *NTracks()  const { return &fCommonMem->fNTracks; }
   GPUhd() AliHLTTPCCATrack *Tracks() const { return  fTracks; }
@@ -237,15 +236,15 @@ class AliHLTTPCCATracker
   
   GPUhd() AliHLTTPCCASliceOutput** Output() const { return fOutput; }
   
-  GPUh() commonMemoryStruct *CommonMemory() {return(fCommonMem); }
+  GPUh() commonMemoryStruct *CommonMemory() const {return(fCommonMem); }
   GPUh() static  size_t CommonMemorySize() { return(sizeof(AliHLTTPCCATracker::commonMemoryStruct)); }
-  GPUh() char* &HitMemory() {return(fHitMemory); }
+  GPUh() char* HitMemory() const {return(fHitMemory); }
   GPUh() size_t HitMemorySize() const {return(fHitMemorySize); }
-  GPUh() char* &TrackletMemory() {return(fTrackletMemory); }
+  GPUh() char* TrackletMemory() {return(fTrackletMemory); }
   GPUh() size_t TrackletMemorySize() const {return(fTrackletMemorySize); }
-  GPUh() char* &TrackMemory() {return(fTrackMemory); }
+  GPUh() char* TrackMemory() {return(fTrackMemory); }
   GPUh() size_t TrackMemorySize() const {return(fTrackMemorySize); }
-  GPUhd() AliHLTTPCCARow* SliceDataRows() {return(fData.Rows()); }
+  GPUhd() AliHLTTPCCARow* SliceDataRows() const {return(fData.Rows()); }
   
   GPUhd() uint3* RowStartHitCountOffset() const {return(fRowStartHitCountOffset);}
   GPUhd() AliHLTTPCCATrackletConstructor::AliHLTTPCCAGPUTempMemory* GPUTrackletTemp() const {return(fGPUTrackletTemp);}
index 491c9cf20878c3a4e6100a7c2d386e8e8410d5c5..3e881a7dc1cf3047b99c42bec97d9498012e0614 100644 (file)
@@ -74,6 +74,7 @@ int AliHLTTPCCATrackerFramework::SetGPUTracker(bool enable)
 
 GPUhd() void AliHLTTPCCATrackerFramework::SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val)
 {
+       //Set Output Control Pointers
        fOutputControl = val;
        fGPUTracker.SetOutputControl(val);
        for (int i = 0;i < fgkNSlices;i++)
index 24723ba2c14d1d18403ba168324ee90af8f7263b..553d4df5c0f409dfb758289417fa8e1bab137bc9 100644 (file)
@@ -206,7 +206,11 @@ GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
        if (r.fStartRow < r.fFirstRow) r.fFirstRow = r.fStartRow;
        tracklet.SetFirstRow( r.fFirstRow );
     tracklet.SetLastRow( r.fLastRow );
+#ifdef HLTCA_GPUCODE
     tracklet.SetParam( tParam.fParam );
+#else
+    tracklet.SetParam( tParam.GetParam() );
+#endif
     int w = ( r.fNHits << 16 ) + r.fItr;
     for ( int iRow = r.fFirstRow; iRow <= r.fLastRow; iRow++ ) {
 #ifdef EXTERN_ROW_HITS