]>
Commit | Line | Data |
---|---|---|
d3821846 | 1 | #define __OPENCL__ |
2 | #define RADEON | |
3 | ||
4 | //Disable assertions since they produce errors in GPU Code | |
5 | #ifdef assert | |
6 | #undef assert | |
7 | #endif | |
8 | #define assert(param) | |
9 | ||
10 | #include "AliHLTTPCCATrackParam.cxx" | |
11 | #include "AliHLTTPCCATrack.cxx" | |
12 | ||
13 | #include "AliHLTTPCCAHitArea.cxx" | |
14 | #include "AliHLTTPCCAGrid.cxx" | |
15 | #include "AliHLTTPCCARow.cxx" | |
16 | #include "AliHLTTPCCAParam.cxx" | |
17 | #include "AliHLTTPCCATracker.cxx" | |
18 | ||
19 | #include "AliHLTTPCCATrackletSelector.cxx" | |
20 | #include "AliHLTTPCCANeighboursFinder.cxx" | |
21 | #include "AliHLTTPCCANeighboursCleaner.cxx" | |
22 | #include "AliHLTTPCCAStartHitsFinder.cxx" | |
23 | #include "AliHLTTPCCAStartHitsSorter.cxx" | |
24 | #include "AliHLTTPCCATrackletConstructor.cxx" | |
25 | ||
26 | __kernel void PreInitRowBlocks(__global char* gpu_mem, GPUconstant() void* pTrackerTmp, int iSlice) | |
27 | { | |
28 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) &pTracker = (( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp)[iSlice]; | |
29 | if (gpu_mem != pTracker.GPUParametersConst()->fGPUMem) return; | |
30 | ||
31 | //Initialize GPU RowBlocks and HitWeights | |
32 | const int nSliceDataHits = pTracker.Data().NumberOfHitsPlusAlign(); | |
33 | __global int4* SliceDataHitWeights4 = (__global int4*) pTracker.Data().HitWeights(); | |
34 | ||
35 | const int stride = get_global_size(0); | |
36 | int4 i0; | |
37 | i0.x = i0.y = i0.z = i0.w = 0; | |
38 | for (int i = get_global_id(0);i < nSliceDataHits * sizeof(int) / sizeof(int4);i += stride) | |
39 | SliceDataHitWeights4[i] = i0; | |
40 | } | |
41 | ||
42 | GPUg() void AliHLTTPCCAProcess_AliHLTTPCCANeighboursFinder(__global char* gpu_mem, GPUconstant() void* pTrackerTmp, int iSlice) | |
43 | { | |
44 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) &pTracker = (( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp)[iSlice]; | |
45 | if (gpu_mem != pTracker.GPUParametersConst()->fGPUMem) return; | |
46 | GPUshared() typename AliHLTTPCCANeighboursFinder::MEM_LOCAL(AliHLTTPCCASharedMemory) smem; | |
47 | ||
48 | for( int iSync=0; iSync<=AliHLTTPCCANeighboursFinder::NThreadSyncPoints(); iSync++){ | |
49 | GPUsync(); | |
50 | AliHLTTPCCANeighboursFinder::Thread( get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), iSync, smem, pTracker ); | |
51 | } | |
52 | } | |
53 | ||
54 | GPUg() void AliHLTTPCCAProcess_AliHLTTPCCANeighboursCleaner(__global char* gpu_mem, GPUconstant() void* pTrackerTmp, int iSlice) | |
55 | { | |
56 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) &pTracker = (( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp)[iSlice]; | |
57 | if (gpu_mem != pTracker.GPUParametersConst()->fGPUMem) return; | |
58 | GPUshared() typename AliHLTTPCCANeighboursCleaner::MEM_LOCAL(AliHLTTPCCASharedMemory) smem; | |
59 | ||
60 | for( int iSync=0; iSync<=AliHLTTPCCANeighboursCleaner::NThreadSyncPoints(); iSync++){ | |
61 | GPUsync(); | |
62 | AliHLTTPCCANeighboursCleaner::Thread( get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), iSync, smem, pTracker ); | |
63 | } | |
64 | } | |
65 | ||
66 | GPUg() void AliHLTTPCCAProcess_AliHLTTPCCAStartHitsFinder(__global char* gpu_mem, GPUconstant() void* pTrackerTmp, int iSlice) | |
67 | { | |
68 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) &pTracker = (( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp)[iSlice]; | |
69 | if (gpu_mem != pTracker.GPUParametersConst()->fGPUMem) return; | |
70 | GPUshared() typename AliHLTTPCCAStartHitsFinder::MEM_LOCAL(AliHLTTPCCASharedMemory) smem; | |
71 | ||
72 | for( int iSync=0; iSync<=AliHLTTPCCAStartHitsFinder::NThreadSyncPoints(); iSync++){ | |
73 | GPUsync(); | |
74 | AliHLTTPCCAStartHitsFinder::Thread( get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), iSync, smem, pTracker ); | |
75 | } | |
76 | } | |
77 | ||
78 | GPUg() void AliHLTTPCCAProcess_AliHLTTPCCAStartHitsSorter(__global char* gpu_mem, GPUconstant() void* pTrackerTmp, int iSlice) | |
79 | { | |
80 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) &pTracker = (( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp)[iSlice]; | |
81 | if (gpu_mem != pTracker.GPUParametersConst()->fGPUMem) return; | |
82 | GPUshared() typename AliHLTTPCCAStartHitsSorter::MEM_LOCAL(AliHLTTPCCASharedMemory) smem; | |
83 | ||
84 | for( int iSync=0; iSync<=AliHLTTPCCAStartHitsSorter::NThreadSyncPoints(); iSync++){ | |
85 | GPUsync(); | |
86 | AliHLTTPCCAStartHitsSorter::Thread( get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), iSync, smem, pTracker ); | |
87 | } | |
88 | } | |
89 | ||
90 | GPUg() void AliHLTTPCCAProcessMulti_AliHLTTPCCATrackletSelector(__global char* gpu_mem, GPUconstant() void* pTrackerTmp, int firstSlice, int nSliceCount) | |
91 | { | |
92 | const int iSlice = nSliceCount * (get_group_id(0) + (get_num_groups(0) % nSliceCount != 0 && nSliceCount * (get_group_id(0) + 1) % get_num_groups(0) != 0)) / get_num_groups(0); | |
93 | const int nSliceBlockOffset = get_num_groups(0) * iSlice / nSliceCount; | |
94 | const int sliceBlockId = get_group_id(0) - nSliceBlockOffset; | |
95 | const int sliceGridDim = get_num_groups(0) * (iSlice + 1) / nSliceCount - get_num_groups(0) * (iSlice) / nSliceCount; | |
96 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) &pTracker = (( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp)[firstSlice + iSlice]; | |
97 | if (gpu_mem != pTracker.GPUParametersConst()->fGPUMem) return; | |
98 | GPUshared() typename AliHLTTPCCATrackletSelector::MEM_LOCAL(AliHLTTPCCASharedMemory) smem; | |
99 | ||
100 | for( int iSync=0; iSync<=AliHLTTPCCATrackletSelector::NThreadSyncPoints(); iSync++){ | |
101 | GPUsync(); | |
102 | AliHLTTPCCATrackletSelector::Thread( sliceGridDim, get_local_size(0), sliceBlockId, get_local_id(0), iSync, smem, pTracker ); | |
103 | } | |
104 | } | |
105 | ||
106 | GPUg() void AliHLTTPCCATrackletConstructorGPU(__global char* gpu_mem, GPUconstant() void* pTrackerTmp) | |
107 | { | |
108 | //GPU Wrapper for AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU | |
109 | GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) *pTracker = ( GPUconstant() MEM_CONSTANT(AliHLTTPCCATracker) * ) pTrackerTmp ; | |
110 | if (gpu_mem != pTracker[0].GPUParametersConst()->fGPUMem) return; | |
111 | GPUshared() AliHLTTPCCATrackletConstructor::MEM_LOCAL(AliHLTTPCCASharedMemory) sMem; | |
112 | AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU(pTracker, sMem); | |
113 | } |