HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h

   1 //-*- Mode: C++ -*-
   2 // ************************************************************************
   3 // This file is property of and copyright by the ALICE HLT Project        *
   4 // ALICE Experiment at CERN, All rights reserved.                         *
   5 // See cxx source for full Copyright notice                               *
   6 //                                                                        *
   7 //*************************************************************************
   8
   9 #ifndef ALIHLTTPCCAPROCESS_H
  10 #define ALIHLTTPCCAPROCESS_H
  11
  12
  13 /**
  14  * Definitions needed for AliHLTTPCCATracker
  15  *
  16  */
  17
  18 #include "AliHLTTPCCADef.h"
  19 #include "AliHLTTPCCATrackParam.h"
  20
  21 class AliHLTTPCCATracker;
  22
  23 #if defined(HLTCA_GPUCODE)
  24
  25 template<class TProcess>
  26 GPUg() void AliHLTTPCCAProcess()
  27 {
  28   AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker );
  29
  30   GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem;
  31
  32   TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, 0, smem, tracker  );
  33
  34 #define GPUPROCESS(iSync) \
  35   if( TProcess::NThreadSyncPoints()>=iSync ){ \
  36     GPUsync(); \
  37     TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker  ); \
  38   }
  39
  40   GPUPROCESS( 1 )
  41   GPUPROCESS( 2 )
  42   GPUPROCESS( 3 )
  43
  44   //for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){
  45   //__syncthreads();
  46   //TProcess::ThreadGPU( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker  );
  47   //}
  48
  49 #undef GPUPROCESS
  50 }
  51
  52 #else
  53
  54 template<class TProcess>
  55 GPUg() void AliHLTTPCCAProcess( int nBlocks, int nThreads, AliHLTTPCCATracker &tracker )
  56 {
  57   for ( int iB = 0; iB < nBlocks; iB++ ) {
  58     typename TProcess::AliHLTTPCCASharedMemory smem;
  59     for ( int iS = 0; iS <= TProcess::NThreadSyncPoints(); iS++ )
  60       for ( int iT = 0; iT < nThreads; iT++ ) {
  61         TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, tracker  );
  62       }
  63   }
  64 }
  65
  66 #endif
  67
  68
  69
  70 #if defined(HLTCA_GPUCODE)
  71
  72 template<typename TProcess>
  73 GPUg() void AliHLTTPCCAProcess1()
  74 {
  75   AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker );
  76   AliHLTTPCCATrackParam tParam;
  77
  78   GPUshared() typename TProcess::AliHLTTPCCASharedMemory sMem;
  79
  80   typename TProcess::AliHLTTPCCAThreadMemory rMem;
  81
  82   for ( int iSync = 0; iSync <= TProcess::NThreadSyncPoints(); iSync++ ) {
  83     GPUsync();
  84     TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync,
  85                       sMem, rMem, tracker, tParam  );
  86   }
  87 }
  88
  89 #else
  90
  91 template<typename TProcess>
  92 GPUg() void AliHLTTPCCAProcess1( int nBlocks, int nThreads, AliHLTTPCCATracker &tracker )
  93 {
  94   for ( int iB = 0; iB < nBlocks; iB++ ) {
  95     typename TProcess::AliHLTTPCCASharedMemory smem;
  96     typename TProcess::AliHLTTPCCAThreadMemory *rMem = new typename TProcess::AliHLTTPCCAThreadMemory[nThreads];
  97     AliHLTTPCCATrackParam *tParam = new AliHLTTPCCATrackParam[ nThreads ];
  98     for ( int iS = 0; iS <= TProcess::NThreadSyncPoints(); iS++ ) {
  99       for ( int iT = 0; iT < nThreads; iT++ )
 100         TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, rMem[iT], tracker, tParam[iT]  );
 101     }
 102     delete[] rMem;
 103     delete[] tParam;
 104   }
 105 }
 106
 107 #endif
 108
 109 #endif