Please consider the following formatting changes to #14542

GPU/GPUTracking/Base/GPUReconstruction.cxx

@@ -271,9 +271,6 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()
 #endif
     mProcessingSettings->overrideClusterizerFragmentLen = TPC_MAX_FRAGMENT_LEN_GPU;
     param().rec.tpc.nWaysOuter = true;
-    if (param().rec.tpc.looperInterpolationInExtraPass == -1) {
-      param().rec.tpc.looperInterpolationInExtraPass = 0;
-    }
     if (GetProcessingSettings().createO2Output > 1) {
       mProcessingSettings->createO2Output = 1;
     }

GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx

@@ -32,7 +32,6 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
   GPUTPCCompression& GPUrestrict() compressor = processors.tpcCompressor;
   const GPUParam& GPUrestrict() param = processors.param;
 
-  uint8_t lastLeg = 0;
   int32_t myTrack = 0;
   for (uint32_t i = get_global_id(0); i < ioPtrs.nMergedTracks; i += get_global_size(0)) {
     GPUbarrierWarp();
@@ -75,9 +74,6 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
         if ((hit.sector < GPUCA_NSECTORS) ^ (lastSector < GPUCA_NSECTORS)) {
           break;
         }
-        if (lastLeg != hit.leg && track.Mirror()) {
-          break;
-        }
         if (track.Propagate(geo.Row2X(hit.row), param.SectorParam[hit.sector].Alpha)) {
           break;
         }
@@ -93,7 +89,6 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
 
         myTrack = CAMath::AtomicAdd(&compressor.mMemory->nStoredTracks, 1u);
         compressor.mAttachedClusterFirstIndex[myTrack] = trk.FirstClusterRef();
-        lastLeg = hit.leg;
         c.qPtA[myTrack] = qpt;
         c.rowA[myTrack] = hit.row;
         c.sliceA[myTrack] = hit.sector;
@@ -114,12 +109,11 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
           sector -= lastSector;
         }
         c.rowDiffA[cidx] = row;
-        c.sliceLegDiffA[cidx] = (hit.leg == lastLeg ? 0 : compressor.NSECTORS) + sector;
+        c.sliceLegDiffA[cidx] = sector;
         float pad = CAMath::Max(0.f, CAMath::Min((float)geo.NPads(GPUCA_ROW_COUNT - 1), track.LinearY2Pad(hit.sector, track.Y(), geo.PadWidth(hit.row), geo.NPads(hit.row))));
         c.padResA[cidx] = orgCl.padPacked - orgCl.packPad(pad);
         float time = CAMath::Max(0.f, geo.LinearZ2Time(hit.sector, track.Z() + zOffset));
         c.timeResA[cidx] = (orgCl.getTimePacked() - orgCl.packTime(time)) & 0xFFFFFF;
-        lastLeg = hit.leg;
       }
       uint16_t qtot = orgCl.qTot, qmax = orgCl.qMax;
       uint8_t sigmapad = orgCl.sigmaPadPacked, sigmatime = orgCl.sigmaTimePacked;

GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h

@@ -47,7 +47,7 @@ struct GPUMemorySizeScalers {
   double tpcSectorTracksPerHit = 0.02;
   double tpcSectorTrackHitsPerHit = 0.8;
   double tpcSectorTrackHitsPerHitWithRejection = 1.0;
-  double tpcMergedTrackPerSectorTrack = 0.9;
+  double tpcMergedTrackPerSectorTrack = 1.0;
   double tpcMergedTrackHitPerSectorHit = 1.1;
   size_t tpcCompressedUnattachedHitsBase1024[3] = {900, 900, 500}; // No ratio, but integer fraction of 1024 for exact computation
 

GPU/GPUTracking/DataTypes/GPUTPCGMMergedTrackHit.h

@@ -21,7 +21,7 @@ namespace o2::gpu
 {
 struct GPUTPCGMMergedTrackHit {
   uint32_t num;
-  uint8_t sector, row, leg, state;
+  uint8_t sector, row, state;
 
   // NOTE: the lower states must match those from ClusterNative!
   // TODO: take them directly from clusterNative header.

GPU/GPUTracking/Definitions/GPUDefConstantsAndSettings.h

@@ -32,7 +32,7 @@
 
 #define GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(QPTB5) (CAMath::Abs(QPTB5) > 10 ? 10 : (CAMath::Abs(QPTB5) > 5 ? 15 : 29)) // Minimum hits should depend on Pt, low Pt tracks can have few hits. 29 Hits default, 15 for < 200 mev, 10 for < 100 mev
 
-#define GPUCA_MERGER_MAX_TRACK_CLUSTERS 1000          // Maximum number of clusters a track may have after merging
+#define GPUCA_MERGER_MAX_TRACK_CLUSTERS 512          // Maximum number of clusters a track may have after merging
 
 #define GPUCA_MAXN 40                                 // Maximum number of neighbor hits to consider in one row in neightbors finder
 #define GPUCA_MIN_TRACK_PTB5_DEFAULT 0.010f           // Default setting for minimum track Pt at some places (at B=0.5T)

GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h

@@ -80,7 +80,6 @@
     #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20
     #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1
     #define GPUCA_PAR_SORT_BEFORE_FIT 1
-    #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1
     #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1
     #define GPUCA_PAR_DEDX_STORAGE_TYPE uint16_t
     #define GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE half
@@ -143,7 +142,6 @@
     #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20
     #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1
     #define GPUCA_PAR_SORT_BEFORE_FIT 1
-    #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1
     #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1
     #define GPUCA_PAR_DEDX_STORAGE_TYPE uint16_t
     #define GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE half
@@ -206,7 +204,6 @@
     #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20
     #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1
     #define GPUCA_PAR_SORT_BEFORE_FIT 1
-    #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1
     #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1
     #define GPUCA_PAR_DEDX_STORAGE_TYPE uint16_t
     #define GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE half
@@ -261,7 +258,6 @@
     #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20
     #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1
     #define GPUCA_PAR_SORT_BEFORE_FIT 1
-    #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1
     #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1
     #define GPUCA_PAR_COMP_GATHER_KERNEL 4
     #define GPUCA_PAR_COMP_GATHER_MODE 3
@@ -529,9 +525,6 @@
   #ifndef GPUCA_PAR_SORT_BEFORE_FIT
     #define GPUCA_PAR_SORT_BEFORE_FIT 0
   #endif
-  #ifndef GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION
-    #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 0
-  #endif
   #ifndef GPUCA_PAR_COMP_GATHER_KERNEL
     #define GPUCA_PAR_COMP_GATHER_KERNEL 0
   #endif
@@ -566,9 +559,6 @@
   #ifndef GPUCA_PAR_SORT_BEFORE_FIT
     #define GPUCA_PAR_SORT_BEFORE_FIT 0
   #endif
-  #ifndef GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION
-    #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 0
-  #endif
   #ifndef GPUCA_PAR_COMP_GATHER_KERNEL
     #define GPUCA_PAR_COMP_GATHER_KERNEL 0
   #endif

GPU/GPUTracking/Definitions/GPUSettingsList.h

@@ -132,11 +132,11 @@ AddOptionRTC(cfInnerThreshold, uint8_t, 0, "", 0, "Cluster Finder extends cluste
 AddOptionRTC(cfMinSplitNum, uint8_t, 1, "", 0, "Minimum number of split charges in a cluster for the cluster to be marked as split")
 AddOptionRTC(cfNoiseSuppressionEpsilon, uint8_t, 10, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression")
 AddOptionRTC(cfNoiseSuppressionEpsilonRelative, uint8_t, 76, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression, relative as fraction of 255")
-AddOptionRTC(cfEdgeTwoPads, uint8_t, 1, "", 0, "Flag clusters with peak on the 2 pads closes to the sector edge as edge cluster")
+AddOptionRTC(cfEdgeTwoPads, uint8_t, 0, "", 0, "Flag clusters with peak on the 2 pads closes to the sector edge as edge cluster")
 AddOptionRTC(nWays, uint8_t, 3, "", 0, "Do N fit passes in final fit of merger")
 AddOptionRTC(nWaysOuter, int8_t, 0, "", 0, "Store outer param")
 AddOptionRTC(trackFitRejectMode, int8_t, 5, "", 0, "0: no limit on rejection or missed hits, >0: break after n rejected hits, <0: reject at max -n hits")
-AddOptionRTC(rejectIFCLowRadiusCluster, uint8_t, 0, "", 0, "Reject clusters that get the IFC mask error during refit")
+AddOptionRTC(rejectIFCLowRadiusCluster, uint8_t, 1, "", 0, "Reject clusters that get the IFC mask error during refit")
 AddOptionRTC(dEdxTruncLow, uint8_t, 2, "", 0, "Low truncation threshold, fraction of 128")
 AddOptionRTC(dEdxTruncHigh, uint8_t, 77, "", 0, "High truncation threshold, fraction of 128")
 AddOptionRTC(extrapolationTracking, int8_t, 1, "", 0, "Enable Extrapolation Tracking (prolong tracks to adjacent sectors to find short segments)")
@@ -154,7 +154,6 @@ AddOptionRTC(mergerInterpolateErrors, uint8_t, 1, "", 0, "Use interpolation inst
 AddOptionRTC(mergerInterpolateRejectAlsoOnCurrentPosition, uint8_t, 1, "", 0, "When using mergerInterpolateErrors, reject based on chi2 twice computed with interpolated and current track position")
 AddOptionRTC(mergeCE, uint8_t, 1, "", 0, "Merge tracks accross the central electrode")
 AddOptionRTC(retryRefit, int8_t, 1, "", 0, "Retry refit with seeding errors and without cluster rejection when fit fails (=2 means retry in same kernel, =1 for separate kernel")
-AddOptionRTC(looperInterpolationInExtraPass, int8_t, -1, "", 0, "Perform looper interpolation in an extra pass")
 AddOptionRTC(dropSecondaryLegsInOutput, int8_t, 1, "", 0, "Do not store secondary legs of looping track in TrackTPC")
 AddOptionRTC(enablePID, int8_t, 1, "", 0, "Enable PID response")
 AddOptionRTC(PID_useNsigma, int8_t, 1, "", 0, "Use nSigma instead of absolute distance in PID response")

GPU/GPUTracking/Global/GPUChainTracking.cxx

@@ -824,14 +824,18 @@ int32_t GPUChainTracking::RunChainFinalize()
 
   if (GetProcessingSettings().eventDisplay) {
     if (!mDisplayRunning) {
+      GPUInfo("Starting Event Display...");
       if (mEventDisplay->StartDisplay()) {
+        GPUError("Error starting Event Display");
         return (1);
       }
       mDisplayRunning = true;
     } else {
       mEventDisplay->ShowNextEvent();
     }
 
+    mEventDisplay->WaitTillEventShown();
+
     if (GetProcessingSettings().eventDisplay->EnableSendKey()) {
       while (kbhit()) {
         getch();
@@ -863,9 +867,9 @@ int32_t GPUChainTracking::RunChainFinalize()
       return (2);
     }
     GetProcessingSettings().eventDisplay->setDisplayControl(0);
-    GPUInfo("Loading next event");
+    GPUInfo("Loading next event...");
 
-    mEventDisplay->WaitForNextEvent();
+    mEventDisplay->BlockTillNextEvent();
   }
 
   return 0;

GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx

@@ -592,7 +592,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
     return ForwardTPCDigits();
   }
 #ifdef GPUCA_TPC_GEOMETRY_O2
-  int32_t tpcTimeBinCut = mUpdateNewCalibObjects && mNewCalibValues->newTPCTimeBinCut ? mNewCalibValues->tpcTimeBinCut : param().tpcCutTimeBin;
+  [[maybe_unused]] int32_t tpcTimeBinCut = mUpdateNewCalibObjects && mNewCalibValues->newTPCTimeBinCut ? mNewCalibValues->tpcTimeBinCut : param().tpcCutTimeBin;
   mRec->PushNonPersistentMemory(qStr2Tag("TPCCLUST"));
   const auto& threadContext = GetThreadContext();
   const bool doGPU = GetRecoStepsGPU() & RecoStep::TPCClusterFinding;
@@ -861,7 +861,6 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
           int32_t firstHBF = (mIOPtrs.settingsTF && mIOPtrs.settingsTF->hasTfStartOrbit) ? mIOPtrs.settingsTF->tfStartOrbit : ((mIOPtrs.tpcZS->sector[iSector].count[0] && mIOPtrs.tpcZS->sector[iSector].nZSPtr[0][0]) ? o2::raw::RDHUtils::getHeartBeatOrbit(*(const o2::header::RAWDataHeader*)mIOPtrs.tpcZS->sector[iSector].zsPtr[0][0]) : 0);
           uint32_t nBlocks = doGPU ? clusterer.mPmemory->counters.nPagesSubsector : GPUTrackingInOutZS::NENDPOINTS;
 
-          (void)tpcTimeBinCut; // TODO: To be used in decoding kernels
           switch (mCFContext->zsVersion) {
             default:
               GPUFatal("Data with invalid TPC ZS mode (%d) received", mCFContext->zsVersion);

GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx

@@ -70,7 +70,7 @@ int32_t GPUChainTracking::DoProfile()
   fwrite(&bmpFH, 1, sizeof(bmpFH), fp2);
   fwrite(&bmpIH, 1, sizeof(bmpIH), fp2);
 
-  int32_t nEmptySync = 0;
+  [[maybe_unused]] int32_t nEmptySync = 0;
   for (uint32_t i = 0; i < bmpheight * ConstructorBlockCount() * ConstructorThreadCount(); i += ConstructorBlockCount() * ConstructorThreadCount()) {
     int32_t fEmpty = 1;
     for (uint32_t j = 0; j < ConstructorBlockCount() * ConstructorThreadCount(); j++) {
@@ -103,7 +103,6 @@ int32_t GPUChainTracking::DoProfile()
     } else {
       nEmptySync = 0;
     }
-    (void)nEmptySync;
     // if (nEmptySync == GPUCA_SCHED_ROW_STEP + 2) break;
   }
 

GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx

@@ -224,9 +224,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
   if (param().rec.tpc.retryRefit == 1) {
     runKernel<GPUTPCGMMergerTrackFit>(GetGridAuto(0), -1);
   }
-  if (param().rec.tpc.looperInterpolationInExtraPass == -1 ? mRec->getGPUParameters(doGPU).par_MERGER_SPLIT_LOOP_INTERPOLATION : param().rec.tpc.looperInterpolationInExtraPass) {
-    runKernel<GPUTPCGMMergerFollowLoopers>(GetGridAuto(0));
-  }
+  runKernel<GPUTPCGMMergerFollowLoopers>(GetGridAuto(0));
 
   DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingRefit, Merger, &GPUTPCGMMerger::DumpRefit, *mDebugFile);
   runKernel<GPUTPCGMMergerFinalize, 0>(GetGridAuto(0, deviceType));

GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx

@@ -28,8 +28,7 @@ int32_t GPUChainTracking::RunRefit()
   GPUTrackingRefitProcessor& Refit = processors()->trackingRefit;
   GPUTrackingRefitProcessor& RefitShadow = doGPU ? processorsShadow()->trackingRefit : Refit;
 
-  const auto& threadContext = GetThreadContext();
-  (void)threadContext;
+  [[maybe_unused]] const auto& threadContext = GetThreadContext();
   SetupGPUProcessor(&Refit, false);
   RefitShadow.SetPtrsFromGPUConstantMem(processorsShadow(), doGPU ? &processorsDevice()->param : nullptr);
   RefitShadow.SetPropagator(doGPU ? processorsShadow()->calibObjects.o2Propagator : GetO2Propagator());

GPU/GPUTracking/Interface/GPUO2InterfaceDisplay.cxx

@@ -46,7 +46,8 @@ int32_t GPUO2InterfaceDisplay::startDisplay()
   if (retVal) {
     return retVal;
   }
-  mDisplay->WaitForNextEvent();
+  mDisplay->WaitTillEventShown();
+  mDisplay->BlockTillNextEvent();
   return 0;
 }
 
@@ -59,14 +60,15 @@ int32_t GPUO2InterfaceDisplay::show(const GPUTrackingInOutPointers* ptrs)
     ptrs = tmpPtr.get();
   }
   mDisplay->ShowNextEvent(ptrs);
+  mDisplay->WaitTillEventShown();
   do {
     usleep(10000);
   } while (mFrontend->getDisplayControl() == 0);
   if (mFrontend->getDisplayControl() == 2) {
     return 1;
   }
   mFrontend->setDisplayControl(0);
-  mDisplay->WaitForNextEvent();
+  mDisplay->BlockTillNextEvent();
   return 0;
 }
 

GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h

@@ -45,13 +45,20 @@ class GPUTPCGMMergedTrack
   GPUd() bool Looper() const { return mFlags & 0x02; }
   GPUd() bool CSide() const { return mFlags & 0x04; }
   GPUd() bool CCE() const { return mFlags & 0x08; }
-  GPUd() bool MergedLooper() const { return mFlags & 0x10; }
+  GPUd() bool MergedLooperUnconnected() const { return mFlags & 0x10; }
+  GPUd() bool MergedLooperConnected() const { return mFlags & 0x20; }
+  GPUd() bool MergedLooper() const { return mFlags & 0x30; }
+  GPUd() int32_t PrevSegment() const { return mPrevSegment; }
+  GPUd() uint8_t Leg() const { return mLeg; }
+  GPUd() uint8_t Flags() const { return mFlags; }
 
   GPUd() void SetNClusters(int32_t v) { mNClusters = v; }
   GPUd() void SetNClustersFitted(int32_t v) { mNClustersFitted = v; }
   GPUd() void SetFirstClusterRef(int32_t v) { mFirstClusterRef = v; }
   GPUd() void SetParam(const GPUTPCGMTrackParam& v) { mParam = v; }
   GPUd() void SetAlpha(float v) { mAlpha = v; }
+  GPUd() void SetPrevSegment(int32_t v) { mPrevSegment = v; }
+  GPUd() void SetLeg(uint8_t v) { mLeg = v; }
   GPUd() void SetOK(bool v)
   {
     if (v) {
@@ -84,18 +91,23 @@ class GPUTPCGMMergedTrack
       mFlags &= 0xF7;
     }
   }
-  GPUd() void SetMergedLooper(bool v)
+  GPUd() void SetMergedLooperUnconnected(bool v)
   {
     if (v) {
       mFlags |= 0x10;
     } else {
       mFlags &= 0xEF;
     }
   }
+  GPUd() void SetMergedLooperConnected(bool v)
+  {
+    if (v) {
+      mFlags |= 0x20;
+    } else {
+      mFlags &= 0xDF;
+    }
+  }
   GPUd() void SetFlags(uint8_t v) { mFlags = v; }
-  GPUd() void SetLegs(uint8_t v) { mLegs = v; }
-  GPUd() uint8_t Legs() const { return mLegs; }
-  GPUd() uint8_t Flags() const { return mFlags; }
 
   GPUd() const gputpcgmmergertypes::GPUTPCOuterParam& OuterParam() const { return mOuterParam; }
   GPUd() gputpcgmmergertypes::GPUTPCOuterParam& OuterParam() { return mOuterParam; }
@@ -106,11 +118,11 @@ class GPUTPCGMMergedTrack
 
   float mAlpha;              //* alpha angle
   uint32_t mFirstClusterRef; //* index of the first track cluster in corresponding cluster arrays
-  // TODO: Change to 8 bit
-  uint32_t mNClusters;       //* number of track clusters
-  uint32_t mNClustersFitted; //* number of clusters used in fit
+  int32_t mPrevSegment;      //* next segment in case of looping track
+  uint16_t mNClusters;       //* number of track clusters
+  uint16_t mNClustersFitted; //* number of clusters used in fit
   uint8_t mFlags;
-  uint8_t mLegs;
+  uint8_t mLeg;
 
 #if !defined(GPUCA_STANDALONE)
   ClassDefNV(GPUTPCGMMergedTrack, 0);