ITS::gpu: Update track selection logics to the state of the art (#13816)

Author: mconcas

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

@@ -40,6 +40,7 @@ GPUg() void fitTrackSeedsKernel(
   CellSeed* trackSeeds,
   const TrackingFrameInfo** foundTrackingFrameInfo,
   o2::its::TrackITSExt* tracks,
+  const float* minPts,
   const unsigned int nSeeds,
   const float Bz,
   const int startLevel,
@@ -182,6 +183,7 @@ void filterCellNeighboursHandler(std::vector<int>&,
 void trackSeedHandler(CellSeed* trackSeeds,
                       const TrackingFrameInfo** foundTrackingFrameInfo,
                       o2::its::TrackITSExt* tracks,
+                      std::vector<float>& minPtsHost,
                       const unsigned int nSeeds,
                       const float Bz,
                       const int startLevel,

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

@@ -278,6 +278,9 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
     const int minimumLayer{startLevel - 1};
     std::vector<CellSeed> trackSeeds;
     for (int startLayer{mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) {
+      if ((mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) {
+        continue;
+      }
       std::vector<int> lastCellId, updatedCellId;
       std::vector<CellSeed> lastCellSeed, updatedCellSeed;
 
@@ -308,6 +311,7 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
     trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),             // CellSeed* trackSeeds,
                      mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo,
                      mTimeFrameGPU->getDeviceTrackITSExt(),            // o2::its::TrackITSExt* tracks,
+                     mTrkParams[iteration].MinPt,                      // std::vector<float>& minPtsHost,
                      trackSeeds.size(),                                // const size_t nSeeds,
                      mBz,                                              // const float Bz,
                      startLevel,                                       // const int startLevel,

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

@@ -267,6 +267,7 @@ GPUg() void fitTrackSeedsKernel(
   CellSeed* trackSeeds,
   const TrackingFrameInfo** foundTrackingFrameInfo,
   o2::its::TrackITSExt* tracks,
+  const float* minPts,
   const unsigned int nSeeds,
   const float Bz,
   const int startLevel,
@@ -317,7 +318,7 @@ GPUg() void fitTrackSeedsKernel(
                           foundTrackingFrameInfo,   // TrackingFrameInfo** trackingFrameInfo,
                           propagator,               // const o2::base::Propagator* propagator,
                           matCorrType);             // o2::base::PropagatorF::MatCorrType matCorrType
-    if (!fitSuccess) {
+    if (!fitSuccess || temporaryTrack.getPt() < minPts[nLayers - temporaryTrack.getNClusters()]) {
       continue;
     }
     tracks[iCurrentTrackSeedIndex] = temporaryTrack;
@@ -1089,6 +1090,7 @@ void filterCellNeighboursHandler(std::vector<int>& neighHost,
 void trackSeedHandler(CellSeed* trackSeeds,
                       const TrackingFrameInfo** foundTrackingFrameInfo,
                       o2::its::TrackITSExt* tracks,
+                      std::vector<float>& minPtsHost,
                       const unsigned int nSeeds,
                       const float Bz,
                       const int startLevel,
@@ -1099,17 +1101,19 @@ void trackSeedHandler(CellSeed* trackSeeds,
                       const int nBlocks,
                       const int nThreads)
 {
+  thrust::device_vector<float> minPts(minPtsHost);
   gpu::fitTrackSeedsKernel<<<nBlocks, nThreads>>>(
-    trackSeeds,               // CellSeed*
-    foundTrackingFrameInfo,   // TrackingFrameInfo**
-    tracks,                   // TrackITSExt*
-    nSeeds,                   // const unsigned int
-    Bz,                       // const float
-    startLevel,               // const int
-    maxChi2ClusterAttachment, // float
-    maxChi2NDF,               // float
-    propagator,               // const o2::base::Propagator*
-    matCorrType);             // o2::base::PropagatorF::MatCorrType
+    trackSeeds,                           // CellSeed*
+    foundTrackingFrameInfo,               // TrackingFrameInfo**
+    tracks,                               // TrackITSExt*
+    thrust::raw_pointer_cast(&minPts[0]), // const float* minPts,
+    nSeeds,                               // const unsigned int
+    Bz,                                   // const float
+    startLevel,                           // const int
+    maxChi2ClusterAttachment,             // float
+    maxChi2NDF,                           // float
+    propagator,                           // const o2::base::Propagator*
+    matCorrType);                         // o2::base::PropagatorF::MatCorrType
 
   gpuCheckError(cudaPeekAtLastError());
   gpuCheckError(cudaDeviceSynchronize());