Please consider the following formatting changes to #14542

Detectors/TPC/reconstruction/test/testGPUCATracking.cxx

@@ -50,24 +50,23 @@ BOOST_AUTO_TEST_CASE(CATracking_test1)
 {
   GPUO2Interface tracker;
 
-  float solenoidBz = -5.00668; //B-field
-  float refX = 1000.;          //transport tracks to this x after tracking, >500 for disabling
-  bool continuous = false;     //time frame data v.s. triggered events
+  float solenoidBz = -5.00668; // B-field
+  float refX = 1000.;          // transport tracks to this x after tracking, >500 for disabling
+  bool continuous = false;     // time frame data v.s. triggered events
 
   GPUO2InterfaceConfiguration config;
   config.configDeviceBackend.deviceType = GPUDataTypes::DeviceType::CPU;
   config.configDeviceBackend.forceDeviceType = true;
 
-  config.configProcessing.ompThreads = 4;         //4 threads if we run on the CPU, 1 = default, 0 = auto-detect
-  config.configProcessing.runQA = false;          //Run QA after tracking
-  config.configProcessing.eventDisplay = nullptr; //Ptr to event display backend, for running standalone OpenGL event display
+  config.configProcessing.ompThreads = 4;         // 4 threads if we run on the CPU, 1 = default, 0 = auto-detect
+  config.configProcessing.runQA = false;          // Run QA after tracking
+  config.configProcessing.eventDisplay = nullptr; // Ptr to event display backend, for running standalone OpenGL event display
 
   config.configGRP.solenoidBzNominalGPU = solenoidBz;
   config.configGRP.grpContinuousMaxTimeBin = continuous ? GPUSettings::TPC_MAX_TF_TIME_BIN : 0; // Number of timebins in timeframe if continuous, 0 otherwise
 
-  config.configReconstruction.tpc.nWays = 3;               //Should always be 3!
-  config.configReconstruction.tpc.nWaysOuter = true;       //Will create outer param for TRD
-  config.configReconstruction.tpc.searchWindowDZDR = 2.5f; //Should always be 2.5 for looper-finding and/or continuous tracking
+  config.configReconstruction.tpc.nWays = 3;               // Should always be 3!
+  config.configReconstruction.tpc.searchWindowDZDR = 2.5f; // Should always be 2.5 for looper-finding and/or continuous tracking
   config.configReconstruction.tpc.trackReferenceX = refX;
 
   config.configWorkflow.steps.set(GPUDataTypes::RecoStep::TPCConversion, GPUDataTypes::RecoStep::TPCSectorTracking,

GPU/Common/GPUCommonMath.h

@@ -74,6 +74,11 @@ class GPUCommonMath
   GPUhdni() constexpr static float Sqrt(float x);
   GPUd() static float InvSqrt(float x);
   template <class T>
+  GPUdi() constexpr static T Square(T x)
+  {
+    return x * x;
+  }
+  template <class T>
   GPUhd() constexpr static T Abs(T x);
   GPUd() constexpr static float ASin(float x);
   GPUd() constexpr static float ACos(float x);

GPU/GPUTracking/Base/GPUConstantMem.h

@@ -23,7 +23,6 @@
 #include "GPUTPCGMMerger.h"
 #include "GPUTRDTracker.h"
 
-#include "GPUTPCConvert.h"
 #include "GPUTPCCompression.h"
 #include "GPUTPCDecompression.h"
 #include "GPUTPCClusterFinder.h"
@@ -42,7 +41,6 @@ namespace o2::gpu
 struct GPUConstantMem {
   GPUParam param;
   GPUTPCTracker tpcTrackers[GPUCA_NSECTORS];
-  GPUTPCConvert tpcConverter;
   GPUTPCCompression tpcCompressor;
   GPUTPCDecompression tpcDecompressor;
   GPUTPCGMMerger tpcMerger;

GPU/GPUTracking/Base/GPUParam.cxx

@@ -111,7 +111,6 @@ void GPUParam::SetDefaults(float solenoidBz, bool assumeConstantBz)
   par.continuousTracking = false;
   continuousMaxTimeBin = 0;
   tpcCutTimeBin = 0;
-  par.earlyTpcTransform = false;
 }
 
 void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessing* p, const GPURecoStepConfiguration* w, const GPUSettingsRecDynamic* d)
@@ -122,7 +121,6 @@ void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessi
     continuousMaxTimeBin = g->grpContinuousMaxTimeBin == -1 ? GPUSettings::TPC_MAX_TF_TIME_BIN : g->grpContinuousMaxTimeBin;
     tpcCutTimeBin = g->tpcCutTimeBin;
   }
-  par.earlyTpcTransform = rec.tpc.forceEarlyTransform == -1 ? (!par.continuousTracking) : rec.tpc.forceEarlyTransform;
   qptB5Scaler = CAMath::Abs(bzkG) > 0.1f ? CAMath::Abs(bzkG) / 5.006680f : 1.f; // Repeat here, since passing in g is optional
   if (p) {
     UpdateRun3ClusterErrors(p->param.tpcErrorParamY, p->param.tpcErrorParamZ);
@@ -156,9 +154,6 @@ void GPUParam::SetDefaults(const GPUSettingsGRP* g, const GPUSettingsRec* r, con
   SetDefaults(g->solenoidBzNominalGPU, g->constBz);
   if (r) {
     rec = *r;
-    if (rec.fitPropagateBzOnly == -1) {
-      rec.fitPropagateBzOnly = rec.tpc.nWays - 1;
-    }
   }
   UpdateSettings(g, p, w);
 }

GPU/GPUTracking/Base/GPUReconstruction.cxx

@@ -269,11 +269,10 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()
 #ifndef GPUCA_DETERMINISTIC_MODE
     GPUError("WARNING, deterministicGPUReconstruction needs GPUCA_DETERMINISTIC_MODE for being fully deterministic, without only most indeterminism by concurrency is removed, but floating point effects remain!");
 #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 (mProcessingSettings->debugLevel >= 6 && ((mProcessingSettings->debugMask + 1) & mProcessingSettings->debugMask)) {
+      GPUError("WARNING: debugMask %d - debug output might not be deterministic with intermediate steps missing", mProcessingSettings->debugMask);
     }
+    mProcessingSettings->overrideClusterizerFragmentLen = TPC_MAX_FRAGMENT_LEN_GPU;
     if (GetProcessingSettings().createO2Output > 1) {
       mProcessingSettings->createO2Output = 1;
     }
@@ -295,14 +294,15 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()
   if (!(mRecoSteps.stepsGPUMask & GPUDataTypes::RecoStep::TPCMerging)) {
     mProcessingSettings->mergerSortTracks = false;
   }
-
   if (GetProcessingSettings().debugLevel > 3 || !IsGPU() || GetProcessingSettings().deterministicGPUReconstruction) {
     mProcessingSettings->delayedOutput = false;
   }
-
   if (!GetProcessingSettings().rtc.enable) {
     mProcessingSettings->rtc.optConstexpr = false;
   }
+  if (GetProcessingSettings().allSanityChecks) {
+    mProcessingSettings->clusterizerZSSanityCheck = mProcessingSettings->mergerSanityCheck = mProcessingSettings->outputSanityCheck = true;
+  }
 
   mMemoryScalers->scalingFactor = GetProcessingSettings().memoryScalingFactor;
   mMemoryScalers->conservative = GetProcessingSettings().conservativeMemoryEstimate;

GPU/GPUTracking/Base/GPUReconstructionCPU.cxx

@@ -17,7 +17,6 @@
 #include "GPUReconstructionThreading.h"
 #include "GPUChain.h"
 #include "GPUDefParametersRuntime.h"
-#include "GPUTPCClusterData.h"
 #include "GPUTPCGMMergedTrack.h"
 #include "GPUTPCGMMergedTrackHit.h"
 #include "GPUTRDTrackletWord.h"

GPU/GPUTracking/CMakeLists.txt

@@ -174,8 +174,6 @@ set(SRCS_NO_CINT ${SRCS_NO_CINT} display/GPUDisplayInterface.cxx)
 set(SRCS_NO_CINT ${SRCS_NO_CINT}
     Global/GPUChainITS.cxx
     dEdx/GPUdEdx.cxx
-    TPCConvert/GPUTPCConvert.cxx
-    TPCConvert/GPUTPCConvertKernel.cxx
     DataCompression/GPUTPCCompression.cxx
     DataCompression/GPUTPCCompressionTrackModel.cxx
     DataCompression/GPUTPCCompressionKernels.cxx

GPU/GPUTracking/DataCompression/GPUTPCClusterRejection.h

@@ -23,7 +23,7 @@ struct GPUTPCClusterRejection {
   template <bool C, class T = void, class S = void>
   static constexpr inline bool GetProtectionStatus(int32_t attach, bool& physics, bool& protect, T* counts = nullptr, S* mev200 = nullptr)
   {
-    (void)counts; // Avoid incorrect -Wunused-but-set-parameter warning
+    (void)counts; // FIXME: Avoid incorrect -Wunused-but-set-parameter warning
     (void)mev200;
     if (attach == 0) {
       return false;

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/GPUDataTypes.h

@@ -101,7 +101,6 @@ class GPUTPCTrack;
 class GPUTPCHitId;
 class GPUTPCGMMergedTrack;
 struct GPUTPCGMMergedTrackHit;
-struct GPUTPCGMMergedTrackHitXYZ;
 class GPUTRDTrackletWord;
 class GPUTRDSpacePoint;
 struct GPUTPCMCInfo;
@@ -237,7 +236,6 @@ struct GPUTrackingInOutPointers {
   const GPUTPCGMMergedTrack* mergedTracks = nullptr;
   uint32_t nMergedTracks = 0;
   const GPUTPCGMMergedTrackHit* mergedTrackHits = nullptr;
-  const GPUTPCGMMergedTrackHitXYZ* mergedTrackHitsXYZ = nullptr;
   uint32_t nMergedTrackHits = 0;
   const uint32_t* mergedTrackHitAttachment = nullptr;
   const uint8_t* mergedTrackHitStates = nullptr;

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.
@@ -38,11 +38,6 @@ struct GPUTPCGMMergedTrackHit {
                   flagHighIncl = 0x80 };
 };
 
-struct GPUTPCGMMergedTrackHitXYZ {
-  float x, y, z;
-  uint16_t amp;
-};
-
 } // namespace o2::gpu
 
 #endif

GPU/GPUTracking/Definitions/GPUDef.h

@@ -60,9 +60,6 @@
   #ifdef CADEBUG
     #undef CADEBUG
   #endif
-  #ifdef GPUCA_CADEBUG_ENABLED
-    #undef GPUCA_CADEBUG_ENABLED
-  #endif
   #if GPUCA_CADEBUG == 1 && !defined(GPUCA_GPUCODE)
     #define CADEBUG(...) __VA_ARGS__
     #define CADEBUG2(cmd, ...) {__VA_ARGS__; cmd;}

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 1024          // 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

@@ -57,9 +57,9 @@
     #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
     #define GPUCA_LB_GPUTPCGMMergerCollect 512
     #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -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
@@ -120,9 +119,9 @@
     #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
     #define GPUCA_LB_GPUTPCGMMergerCollect 512
     #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -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
@@ -183,9 +181,9 @@
     #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
     #define GPUCA_LB_GPUTPCGMMergerCollect 256, 2
     #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -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
@@ -246,9 +243,9 @@
     #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
     #define GPUCA_LB_GPUTPCGMMergerCollect 128, 2
     #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
     #define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -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
@@ -309,9 +305,6 @@
   #ifndef GPUCA_LB_GPUTRDTrackerKernels_o2Version
     #define GPUCA_LB_GPUTRDTrackerKernels_o2Version 512
   #endif
-  #ifndef GPUCA_LB_GPUTPCConvertKernel
-    #define GPUCA_LB_GPUTPCConvertKernel 256
-  #endif
   #ifndef GPUCA_LB_GPUTPCCompressionKernels_step0attached
     #define GPUCA_LB_GPUTPCCompressionKernels_step0attached 256
   #endif
@@ -405,14 +398,14 @@
   #ifndef GPUCA_LB_GPUTPCGMMergerSortTracksPrepare
     #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
   #endif
-  #ifndef GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
+  #ifndef GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
   #endif
-  #ifndef GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
+  #ifndef GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
   #endif
-  #ifndef GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2
-    #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
+  #ifndef GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2
+    #define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
   #endif
   #ifndef GPUCA_LB_GPUTPCGMMergerFinalize_step0
     #define GPUCA_LB_GPUTPCGMMergerFinalize_step0 256
@@ -529,9 +522,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 +556,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