]>
Commit | Line | Data |
---|---|---|
00d07bcd | 1 | //-*- Mode: C++ -*- |
ce565086 | 2 | // ************************************************************************ |
fbb9b71b | 3 | // This file is property of and copyright by the ALICE HLT Project * |
ce565086 | 4 | // ALICE Experiment at CERN, All rights reserved. * |
5 | // See cxx source for full Copyright notice * | |
6 | // * | |
7 | //************************************************************************* | |
00d07bcd | 8 | |
9 | #ifndef ALIHLTTPCCAPROCESS_H | |
10 | #define ALIHLTTPCCAPROCESS_H | |
11 | ||
12 | ||
13 | /** | |
14 | * Definitions needed for AliHLTTPCCATracker | |
15 | * | |
16 | */ | |
17 | ||
18 | #include "AliHLTTPCCADef.h" | |
4687b8fc | 19 | #include "AliHLTTPCCATrackParam.h" |
00d07bcd | 20 | |
21 | class AliHLTTPCCATracker; | |
22 | ||
23 | #if defined(HLTCA_GPUCODE) | |
24 | ||
25 | template<class TProcess> | |
b22af1bf | 26 | GPUg() void AliHLTTPCCAProcess(int iSlice) |
00d07bcd | 27 | { |
b22af1bf | 28 | AliHLTTPCCATracker &tracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker )[iSlice]; |
00d07bcd | 29 | GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem; |
30 | ||
b22af1bf | 31 | for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){ |
32 | __syncthreads(); | |
33 | TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker ); | |
00d07bcd | 34 | } |
b22af1bf | 35 | } |
00d07bcd | 36 | |
f0bada7f | 37 | template <class TProcess> |
38 | GPUg() void AliHLTTPCCAProcessMultiA(int firstSlice, int nSliceCount, int nVirtualBlocks) | |
39 | { | |
40 | if (blockIdx.x >= nSliceCount) return; | |
41 | AliHLTTPCCATracker &tracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker )[firstSlice + blockIdx.x]; | |
42 | ||
43 | GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem; | |
44 | ||
45 | for (int i = 0;i < nVirtualBlocks;i++) | |
46 | { | |
47 | for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){ | |
48 | __syncthreads(); | |
49 | TProcess::Thread( nVirtualBlocks, blockDim.x, i, threadIdx.x, iSync, smem, tracker ); | |
50 | } | |
51 | } | |
52 | } | |
53 | ||
b22af1bf | 54 | template<class TProcess> |
55 | GPUg() void AliHLTTPCCAProcessMulti(int firstSlice, int nSliceCount) | |
56 | { | |
57 | const int iSlice = nSliceCount * (blockIdx.x + (gridDim.x % nSliceCount != 0 && nSliceCount * (blockIdx.x + 1) % gridDim.x != 0)) / gridDim.x; | |
58 | const int nSliceBlockOffset = gridDim.x * iSlice / nSliceCount; | |
59 | const int sliceBlockId = blockIdx.x - nSliceBlockOffset; | |
60 | const int sliceGridDim = gridDim.x * (iSlice + 1) / nSliceCount - gridDim.x * (iSlice) / nSliceCount; | |
61 | AliHLTTPCCATracker &tracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker )[firstSlice + iSlice]; | |
62 | GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem; | |
fbb9b71b | 63 | |
b22af1bf | 64 | for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){ |
65 | __syncthreads(); | |
66 | TProcess::Thread( sliceGridDim, blockDim.x, sliceBlockId, threadIdx.x, iSync, smem, tracker ); | |
67 | } | |
00d07bcd | 68 | } |
69 | ||
70 | #else | |
71 | ||
72 | template<class TProcess> | |
fbb9b71b | 73 | GPUg() void AliHLTTPCCAProcess( int nBlocks, int nThreads, AliHLTTPCCATracker &tracker ) |
74 | { | |
75 | for ( int iB = 0; iB < nBlocks; iB++ ) { | |
00d07bcd | 76 | typename TProcess::AliHLTTPCCASharedMemory smem; |
fbb9b71b | 77 | for ( int iS = 0; iS <= TProcess::NThreadSyncPoints(); iS++ ) |
78 | for ( int iT = 0; iT < nThreads; iT++ ) { | |
79 | TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, tracker ); | |
693d2443 | 80 | } |
00d07bcd | 81 | } |
82 | } | |
83 | ||
31649d4b | 84 | #endif //HLTCA_GPUCODE |
00d07bcd | 85 | |
86 | ||
87 | ||
88 | #if defined(HLTCA_GPUCODE) | |
89 | ||
90 | template<typename TProcess> | |
91 | GPUg() void AliHLTTPCCAProcess1() | |
92 | { | |
7be9b0d7 | 93 | AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker ); |
fbb9b71b | 94 | AliHLTTPCCATrackParam tParam; |
95 | ||
96 | GPUshared() typename TProcess::AliHLTTPCCASharedMemory sMem; | |
97 | ||
4687b8fc | 98 | typename TProcess::AliHLTTPCCAThreadMemory rMem; |
00d07bcd | 99 | |
fbb9b71b | 100 | for ( int iSync = 0; iSync <= TProcess::NThreadSyncPoints(); iSync++ ) { |
101 | GPUsync(); | |
102 | TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, | |
103 | sMem, rMem, tracker, tParam ); | |
104 | } | |
00d07bcd | 105 | } |
106 | ||
107 | #else | |
108 | ||
109 | template<typename TProcess> | |
fbb9b71b | 110 | GPUg() void AliHLTTPCCAProcess1( int nBlocks, int nThreads, AliHLTTPCCATracker &tracker ) |
00d07bcd | 111 | { |
fbb9b71b | 112 | for ( int iB = 0; iB < nBlocks; iB++ ) { |
00d07bcd | 113 | typename TProcess::AliHLTTPCCASharedMemory smem; |
e1f2d1c3 | 114 | typename TProcess::AliHLTTPCCAThreadMemory *rMem = new typename TProcess::AliHLTTPCCAThreadMemory[nThreads]; |
115 | AliHLTTPCCATrackParam *tParam = new AliHLTTPCCATrackParam[ nThreads ]; | |
fbb9b71b | 116 | for ( int iS = 0; iS <= TProcess::NThreadSyncPoints(); iS++ ) { |
117 | for ( int iT = 0; iT < nThreads; iT++ ) | |
118 | TProcess::Thread( nBlocks, nThreads, iB, iT, iS, smem, rMem[iT], tracker, tParam[iT] ); | |
00d07bcd | 119 | } |
e1f2d1c3 | 120 | delete[] rMem; |
121 | delete[] tParam; | |
00d07bcd | 122 | } |
123 | } | |
124 | ||
31649d4b | 125 | #endif //HLTCA_GPUCODE |
00d07bcd | 126 | |
31649d4b | 127 | #endif //ALIHLTTPCCAPROCESS_H |