fix6

Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>

Author: f3sch

Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h

@@ -197,28 +197,6 @@ void processNeighboursHandler(const int startLevel,
                               const o2::base::PropagatorF::MatCorrType matCorrType,
                               o2::its::ExternalAllocator* alloc);
 
-void processNeighboursHandler7(const int startLevel,
-                               const int defaultCellTopologyId,
-                               CellSeed** allCellSeeds,
-                               CellSeed* currentCellSeeds,
-                               const int* currentCellTopologyIds,
-                               const int* currentCellIds,
-                               std::array<int, TrackingTopology<7>::MaxCells>& nCells,
-                               const unsigned char** usedClusters,
-                               CellNeighbour** neighbours,
-                               int** neighboursDeviceLUTs,
-                               const TrackingFrameInfo** foundTrackingFrameInfo,
-                               bounded_vector<TrackSeed<7>>& seedsHost,
-                               const float bz,
-                               const float MaxChi2ClusterAttachment,
-                               const float maxChi2NDF,
-                               const int maxHoles,
-                               const LayerMask holeLayerMask,
-                               const std::vector<float>& layerxX0Host,
-                               const o2::base::Propagator* propagator,
-                               const o2::base::PropagatorF::MatCorrType matCorrType,
-                               o2::its::ExternalAllocator* alloc);
-
 template <int NLayers>
 void countTrackSeedHandler(TrackSeed<NLayers>* trackSeeds,
                            const TrackingFrameInfo** foundTrackingFrameInfo,

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx

@@ -313,29 +313,27 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
           mTimeFrameGPU->getNCells()[startCellTopologyId] == 0) {
         continue;
       }
-      if constexpr (NLayers == 7) {
-        processNeighboursHandler7(startLevel,
-                                  startCellTopologyId,
-                                  mTimeFrameGPU->getDeviceArrayCells(),
-                                  mTimeFrameGPU->getDeviceCells()[startCellTopologyId],
-                                  nullptr,
-                                  nullptr,
-                                  mTimeFrameGPU->getArrayNCells(),
-                                  (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
-                                  mTimeFrameGPU->getDeviceArrayNeighbours(),
-                                  mTimeFrameGPU->getDeviceArrayNeighboursCellLUT(),
-                                  mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
-                                  trackSeeds,
-                                  this->mBz,
-                                  this->mTrkParams[iteration].MaxChi2ClusterAttachment,
-                                  this->mTrkParams[iteration].MaxChi2NDF,
-                                  this->mTrkParams[iteration].MaxHoles,
-                                  this->mTrkParams[iteration].HoleLayerMask,
-                                  this->mTrkParams[iteration].LayerxX0,
-                                  mTimeFrameGPU->getDevicePropagator(),
-                                  this->mTrkParams[iteration].CorrType,
-                                  mTimeFrameGPU->getFrameworkAllocator());
-      }
+      processNeighboursHandler(startLevel,
+                               startCellTopologyId,
+                               mTimeFrameGPU->getDeviceArrayCells(),
+                               mTimeFrameGPU->getDeviceCells()[startCellTopologyId],
+                               nullptr,
+                               nullptr,
+                               mTimeFrameGPU->getArrayNCells(),
+                               (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
+                               mTimeFrameGPU->getDeviceArrayNeighbours(),
+                               mTimeFrameGPU->getDeviceArrayNeighboursCellLUT(),
+                               mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
+                               trackSeeds,
+                               this->mBz,
+                               this->mTrkParams[iteration].MaxChi2ClusterAttachment,
+                               this->mTrkParams[iteration].MaxChi2NDF,
+                               this->mTrkParams[iteration].MaxHoles,
+                               this->mTrkParams[iteration].HoleLayerMask,
+                               this->mTrkParams[iteration].LayerxX0,
+                               mTimeFrameGPU->getDevicePropagator(),
+                               this->mTrkParams[iteration].CorrType,
+                               mTimeFrameGPU->getFrameworkAllocator());
     }
     // fixme: I don't want to move tracks back and forth, but I need a way to use a thrust::allocator that is aware of our managed memory.
     if (trackSeeds.empty()) {

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu

@@ -1028,51 +1028,6 @@ void processNeighboursHandler(const int startLevel,
   alloc->popTagOffStack(Tag);
 }
 
-void processNeighboursHandler7(const int startLevel,
-                               const int defaultCellTopologyId,
-                               CellSeed** allCellSeeds,
-                               CellSeed* currentCellSeeds,
-                               const int* currentCellTopologyIds,
-                               const int* currentCellIds,
-                               std::array<int, TrackingTopology<7>::MaxCells>& nCells,
-                               const unsigned char** usedClusters,
-                               CellNeighbour** neighbours,
-                               int** neighboursDeviceLUTs,
-                               const TrackingFrameInfo** foundTrackingFrameInfo,
-                               bounded_vector<TrackSeed<7>>& seedsHost,
-                               const float bz,
-                               const float maxChi2ClusterAttachment,
-                               const float maxChi2NDF,
-                               const int maxHoles,
-                               const LayerMask holeLayerMask,
-                               const std::vector<float>& layerxX0Host,
-                               const o2::base::Propagator* propagator,
-                               const o2::base::PropagatorF::MatCorrType matCorrType,
-                               o2::its::ExternalAllocator* alloc)
-{
-  processNeighboursHandler<7>(startLevel,
-                              defaultCellTopologyId,
-                              allCellSeeds,
-                              currentCellSeeds,
-                              currentCellTopologyIds,
-                              currentCellIds,
-                              nCells,
-                              usedClusters,
-                              neighbours,
-                              neighboursDeviceLUTs,
-                              foundTrackingFrameInfo,
-                              seedsHost,
-                              bz,
-                              maxChi2ClusterAttachment,
-                              maxChi2NDF,
-                              maxHoles,
-                              holeLayerMask,
-                              layerxX0Host,
-                              propagator,
-                              matCorrType,
-                              alloc);
-}
-
 template <int NLayers>
 void countTrackSeedHandler(TrackSeed<NLayers>* trackSeeds,
                            const TrackingFrameInfo** foundTrackingFrameInfo,