HLT/TPCLib/tracking-ca/cagpu/AliHLTTPCCAGPUTrackerOpenCL.cl

   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 }