Obsolete tracker output removed

[u/mrichter/AliRoot.git] / HLT / TPCLib / tracking-ca / AliHLTTPCCAProcess.h

diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h
index 67556fce507eb5e61bf45805b4e2ce078a376af6..872066fc9e65ed341016d71ccb3239c18893c056 100644 (file)
--- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h
+++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h
@@ -1,6 +1,6 @@
 //-*- Mode: C++ -*-
 // ************************************************************************
-// This file is property of and copyright by the ALICE HLT Project        * 
+// 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                               *
 //                                                                        *
@@ -23,47 +23,48 @@ class AliHLTTPCCATracker;
 #if defined(HLTCA_GPUCODE)
 
 template<class TProcess>
-GPUg() void AliHLTTPCCAProcess()
+GPUg() void AliHLTTPCCAProcess(int iSlice)
 {
-  AliHLTTPCCATracker &tracker = *((AliHLTTPCCATracker*) cTracker);
-
+  AliHLTTPCCATracker &tracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker )[iSlice];
   GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem;
 
-  TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, 0, smem, tracker  );
-
-#define GPUPROCESS(iSync) \
-  if( TProcess::NThreadSyncPoints()>=iSync ){ \
-    GPUsync(); \
-    TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker  ); \
+  for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){
+    __syncthreads();
+    TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker  );
   }
+}
+
+template<class TProcess>
+GPUg() void AliHLTTPCCAProcessMulti(int firstSlice, int nSliceCount)
+{
+  const int iSlice = nSliceCount * (blockIdx.x + (gridDim.x % nSliceCount != 0 && nSliceCount * (blockIdx.x + 1) % gridDim.x != 0)) / gridDim.x;
+  const int nSliceBlockOffset = gridDim.x * iSlice / nSliceCount;
+  const int sliceBlockId = blockIdx.x - nSliceBlockOffset;
+  const int sliceGridDim = gridDim.x * (iSlice + 1) / nSliceCount - gridDim.x * (iSlice) / nSliceCount;
+  AliHLTTPCCATracker &tracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker )[firstSlice + iSlice];
+  GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem;
 
-  GPUPROCESS(1)
-  GPUPROCESS(2)
-  GPUPROCESS(3)
-    
-    //for( Int_t iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){
-    //__syncthreads();
-    //TProcess::ThreadGPU( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker  ); 
-    //}
-    
-#undef GPUPROCESS
+  for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){
+    __syncthreads();
+    TProcess::Thread( sliceGridDim, blockDim.x, sliceBlockId, threadIdx.x, iSync, smem, tracker  );
+  }
 }
 
 #else
 
 template<class TProcess>
-GPUg() void AliHLTTPCCAProcess( Int_t nBlocks, Int_t nThreads, AliHLTTPCCATracker &tracker )
-{  
-  for( Int_t iB=0; iB<nBlocks; iB++ ){
+GPUg() void AliHLTTPCCAProcess( int nBlocks, int nThreads, AliHLTTPCCATracker &tracker )
+{
+  for ( int iB = 0; iB < nBlocks; iB++ ) {
     typename TProcess::AliHLTTPCCASharedMemory smem;
-    for( Int_t iS=0; iS<=TProcess::NThreadSyncPoints(); iS++)
-      for( Int_t iT=0; iT<nThreads; iT++ ){    
-       TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, tracker  );
+    for ( int iS = 0; iS <= TProcess::NThreadSyncPoints(); iS++ )
+      for ( int iT = 0; iT < nThreads; iT++ ) {
+        TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, tracker  );
       }
   }
 }
 
-#endif
+#endif //HLTCA_GPUCODE
 
 
 
@@ -72,36 +73,38 @@ GPUg() void AliHLTTPCCAProcess( Int_t nBlocks, Int_t nThreads, AliHLTTPCCATracke
 template<typename TProcess>
 GPUg() void AliHLTTPCCAProcess1()
 {
-  AliHLTTPCCATracker &tracker = *((AliHLTTPCCATracker*) cTracker);
-  AliHLTTPCCATrackParam tParam; 
-  
-  GPUshared() typename TProcess::AliHLTTPCCASharedMemory sMem;  
-  
+  AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker );
+  AliHLTTPCCATrackParam tParam;
+
+  GPUshared() typename TProcess::AliHLTTPCCASharedMemory sMem;
+
   typename TProcess::AliHLTTPCCAThreadMemory rMem;
 
-  for( Int_t iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){
-    GPUsync(); 
-    TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, 
-                     sMem, rMem, tracker, tParam  ); 
-  }  
+  for ( int iSync = 0; iSync <= TProcess::NThreadSyncPoints(); iSync++ ) {
+    GPUsync();
+    TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync,
+                      sMem, rMem, tracker, tParam  );
+  }
 }
 
 #else
 
 template<typename TProcess>
-GPUg() void AliHLTTPCCAProcess1( Int_t nBlocks, Int_t nThreads, AliHLTTPCCATracker &tracker )
+GPUg() void AliHLTTPCCAProcess1( int nBlocks, int nThreads, AliHLTTPCCATracker &tracker )
 {
-  for( Int_t iB=0; iB<nBlocks; iB++ ){
+  for ( int iB = 0; iB < nBlocks; iB++ ) {
     typename TProcess::AliHLTTPCCASharedMemory smem;
-    typename TProcess::AliHLTTPCCAThreadMemory rMem[nThreads];
-    AliHLTTPCCATrackParam tParam[nThreads];
-    for( Int_t iS=0; iS<=TProcess::NThreadSyncPoints(); iS++){
-      for( Int_t iT=0; iT<nThreads; iT++ )
-       TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, rMem[iT], tracker, tParam[iT]  );
+    typename TProcess::AliHLTTPCCAThreadMemory *rMem = new typename TProcess::AliHLTTPCCAThreadMemory[nThreads];
+    AliHLTTPCCATrackParam *tParam = new AliHLTTPCCATrackParam[ nThreads ];
+    for ( int iS = 0; iS <= TProcess::NThreadSyncPoints(); iS++ ) {
+      for ( int iT = 0; iT < nThreads; iT++ )
+        TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, rMem[iT], tracker, tParam[iT]  );
     }
+    delete[] rMem;
+    delete[] tParam;
   }
 }
 
-#endif
+#endif //HLTCA_GPUCODE
 
-#endif
+#endif //ALIHLTTPCCAPROCESS_H