DO NOT COMMIT

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

Author: f3sch

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

@@ -197,6 +197,28 @@ 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,27 +313,29 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
           mTimeFrameGPU->getNCells()[startCellTopologyId] == 0) {
         continue;
       }
-      processNeighboursHandler<NLayers>(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());
+      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());
+      }
     }
     // 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,6 +1028,51 @@ 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,