GPU: Fix some settings must not be constexpr with RTC since they can change

davidrohr · davidrohr · commit 18b692ecf424 · 2024-09-14T15:27:21.000+02:00
diff --git a/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx b/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx
@@ -148,7 +148,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
     if (clusters.nTracks && clusters.solenoidBz != -1e6f && clusters.solenoidBz != mParam->bzkG) {
       throw std::runtime_error("Configured solenoid Bz does not match value used for track model encoding");
     }
-    if (clusters.nTracks && clusters.maxTimeBin != -1e6 && clusters.maxTimeBin != mParam->par.continuousMaxTimeBin) {
+    if (clusters.nTracks && clusters.maxTimeBin != -1e6 && clusters.maxTimeBin != mParam->continuousMaxTimeBin) {
       throw std::runtime_error("Configured max time bin does not match value used for track model encoding");
     }
     mCTFCoder.setSelectedIRFrames(pc.inputs().get<gsl::span<o2::dataformats::IRFrame>>("selIRFrames"));
@@ -162,7 +162,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
     const float totalT = std::max(mFastTransform->getMaxDriftTime(0), mFastTransform->getMaxDriftTime(GPUCA_NSLICES / 2));
 
     unsigned int offset = 0, lasti = 0;
-    const unsigned int maxTime = (mParam->par.continuousMaxTimeBin + 1) * o2::tpc::ClusterNative::scaleTimePacked - 1;
+    const unsigned int maxTime = (mParam->continuousMaxTimeBin + 1) * o2::tpc::ClusterNative::scaleTimePacked - 1;
 #ifdef WITH_OPENMP
 #pragma omp parallel for firstprivate(offset, lasti) num_threads(mNThreads)
 #endif
diff --git a/GPU/GPUTracking/Base/GPUParam.cxx b/GPU/GPUTracking/Base/GPUParam.cxx
@@ -119,7 +119,7 @@ void GPUParam::SetDefaults(float solenoidBz)
   par.assumeConstantBz = false;
   par.toyMCEventsFlag = false;
   par.continuousTracking = false;
-  par.continuousMaxTimeBin = 0;
+  continuousMaxTimeBin = 0;
   par.debugLevel = 0;
   par.earlyTpcTransform = false;
 }
@@ -131,7 +131,7 @@ void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessi
     par.assumeConstantBz = g->constBz;
     par.toyMCEventsFlag = g->homemadeEvents;
     par.continuousTracking = g->continuousMaxTimeBin != 0;
-    par.continuousMaxTimeBin = g->continuousMaxTimeBin == -1 ? GPUSettings::TPC_MAX_TF_TIME_BIN : g->continuousMaxTimeBin;
+    continuousMaxTimeBin = g->continuousMaxTimeBin == -1 ? GPUSettings::TPC_MAX_TF_TIME_BIN : g->continuousMaxTimeBin;
   }
   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
@@ -140,9 +140,9 @@ void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessi
     UpdateRun3ClusterErrors(p->param.tpcErrorParamY, p->param.tpcErrorParamZ);
   }
   if (w) {
-    par.dodEdx = w->steps.isSet(GPUDataTypes::RecoStep::TPCdEdx);
+    par.dodEdx = dodEdxDownscaled = w->steps.isSet(GPUDataTypes::RecoStep::TPCdEdx);
     if (par.dodEdx && p && p->tpcDownscaledEdx != 0) {
-      par.dodEdx = (rand() % 100) < p->tpcDownscaledEdx;
+      dodEdxDownscaled = (rand() % 100) < p->tpcDownscaledEdx;
     }
   }
 }
diff --git a/GPU/GPUTracking/Base/GPUParam.h b/GPU/GPUTracking/Base/GPUParam.h
@@ -57,6 +57,9 @@ struct GPUParam_t {
   float bzCLight;
   float qptB5Scaler;
 
+  signed char dodEdxDownscaled;
+  int continuousMaxTimeBin;
+
   GPUTPCGeometry tpcGeometry;                       // TPC Geometry
   GPUTPCGMPolynomialField polynomialField;          // Polynomial approx. of magnetic field for TPC GM
   const unsigned int* occupancyMap;                 // Ptr to TPC occupancy map
diff --git a/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx
@@ -28,7 +28,7 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste
   CompressedClusters& GPUrestrict() cmprClusters = decompressor.mInputGPU;
   const GPUParam& GPUrestrict() param = processors.param;
 
-  const unsigned int maxTime = (param.par.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1;
+  const unsigned int maxTime = (param.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1;
 
   for (int i = trackStart + get_global_id(0); i < trackEnd; i += get_global_size(0)) {
     decompressTrack(cmprClusters, param, maxTime, i, decompressor.mAttachedClustersOffsets[i], decompressor);
@@ -85,7 +85,7 @@ GPUdii() void GPUTPCDecompressionKernels::decompressTrack(CompressedClusters& cm
           pad = param.tpcGeometry.NPads(row) * ClusterNative::scalePadPacked - 1;
         }
       }
-      if (param.par.continuousMaxTimeBin > 0 && time >= maxTime) {
+      if (param.continuousMaxTimeBin > 0 && time >= maxTime) {
         if (time >= 0xFFFFFF - 544768) { // Constant 544768 = (2^23 - LHCMAXBUNCHES(3564) * MAXORBITS(256) * scaleTimePacked(64) / BCPERTIMEBIN(8)) / 2)
           time = 0;
         } else {
@@ -152,8 +152,8 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste
         if (t < 0) {
           t = 0;
         }
-        if (processors.param.par.continuousMaxTimeBin > 0 && t > processors.param.par.continuousMaxTimeBin) {
-          t = processors.param.par.continuousMaxTimeBin;
+        if (processors.param.continuousMaxTimeBin > 0 && t > processors.param.continuousMaxTimeBin) {
+          t = processors.param.continuousMaxTimeBin;
         }
         cl.setTime(t);
       }
diff --git a/GPU/GPUTracking/DataCompression/TPCClusterDecompressor.cxx b/GPU/GPUTracking/DataCompression/TPCClusterDecompressor.cxx
@@ -43,7 +43,7 @@ int TPCClusterDecompressor::decompress(const CompressedClusters* clustersCompres
   if (clustersCompressed->nTracks && clustersCompressed->solenoidBz != -1e6f && clustersCompressed->solenoidBz != param.bzkG) {
     throw std::runtime_error("Configured solenoid Bz does not match value used for track model encoding");
   }
-  if (clustersCompressed->nTracks && clustersCompressed->maxTimeBin != -1e6 && clustersCompressed->maxTimeBin != param.par.continuousMaxTimeBin) {
+  if (clustersCompressed->nTracks && clustersCompressed->maxTimeBin != -1e6 && clustersCompressed->maxTimeBin != param.continuousMaxTimeBin) {
     throw std::runtime_error("Configured max time bin does not match value used for track model encoding");
   }
   std::vector<ClusterNative> clusters[NSLICES][GPUCA_ROW_COUNT];
@@ -52,7 +52,7 @@ int TPCClusterDecompressor::decompress(const CompressedClusters* clustersCompres
     (&locks[0][0])[i].clear();
   }
   unsigned int offset = 0, lasti = 0;
-  const unsigned int maxTime = param.par.continuousMaxTimeBin > 0 ? ((param.par.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1) : TPC_MAX_TIME_BIN_TRIGGERED;
+  const unsigned int maxTime = param.continuousMaxTimeBin > 0 ? ((param.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1) : TPC_MAX_TIME_BIN_TRIGGERED;
   GPUCA_OPENMP(parallel for firstprivate(offset, lasti))
   for (unsigned int i = 0; i < clustersCompressed->nTracks; i++) {
     if (i < lasti) {
@@ -99,8 +99,8 @@ int TPCClusterDecompressor::decompress(const CompressedClusters* clustersCompres
           if (t < 0) {
             t = 0;
           }
-          if (param.par.continuousMaxTimeBin > 0 && t > param.par.continuousMaxTimeBin) {
-            t = param.par.continuousMaxTimeBin;
+          if (param.continuousMaxTimeBin > 0 && t > param.continuousMaxTimeBin) {
+            t = param.continuousMaxTimeBin;
           }
           cl.setTime(t);
         }
diff --git a/GPU/GPUTracking/DataCompression/TPCClusterDecompressor.inc b/GPU/GPUTracking/DataCompression/TPCClusterDecompressor.inc
@@ -101,7 +101,7 @@ inline void TPCClusterDecompressor::decompressTrack(const CompressedClusters* cl
           pad = param.tpcGeometry.NPads(row) * ClusterNative::scalePadPacked - 1;
         }
       }
-      if (param.par.continuousMaxTimeBin > 0 && time >= maxTime) {
+      if (param.continuousMaxTimeBin > 0 && time >= maxTime) {
         if (time >= 0xFFFFFF - 544768) { // Constant 544768 = (2^23 - LHCMAXBUNCHES(3564) * MAXORBITS(256) * scaleTimePacked(64) / BCPERTIMEBIN(8)) / 2)
           time = 0;
         } else {
diff --git a/GPU/GPUTracking/DataTypes/GPUTPCClusterOccupancyMap.cxx b/GPU/GPUTracking/DataTypes/GPUTPCClusterOccupancyMap.cxx
@@ -22,7 +22,7 @@ GPUd() unsigned int GPUTPCClusterOccupancyMapBin::getNBins(const GPUParam& param
   if (param.rec.tpc.occupancyMapTimeBins == 0) {
     return 0;
   }
-  unsigned int maxTimeBin = param.par.continuousTracking ? param.par.continuousMaxTimeBin : TPC_MAX_TIME_BIN_TRIGGERED;
+  unsigned int maxTimeBin = param.par.continuousTracking ? param.continuousMaxTimeBin : TPC_MAX_TIME_BIN_TRIGGERED;
   return (maxTimeBin + param.rec.tpc.occupancyMapTimeBins) / param.rec.tpc.occupancyMapTimeBins; // Not -1, since maxTimeBin is allowed
 }
 
diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h
@@ -565,14 +565,13 @@ EndConfig()
 
 // Derrived parameters used in GPUParam
 BeginHiddenConfig(GPUSettingsParam, param)
-AddVariableRTC(dAlpha, float, 0.f)           // angular size
-AddVariableRTC(assumeConstantBz, signed char, 0)    // Assume a constant magnetic field
-AddVariableRTC(toyMCEventsFlag, signed char, 0)     // events were build with home-made event generator
-AddVariableRTC(continuousTracking, signed char, 0)  // Continuous tracking, estimate bz and errors for abs(z) = 125cm during seeding
-AddVariableRTC(dodEdx, signed char, 0)              // Do dEdx computation
-AddVariableRTC(earlyTpcTransform, signed char, 0)   // do Early TPC transformation
-AddVariableRTC(debugLevel, signed char, 0)          // Debug level
-AddVariableRTC(continuousMaxTimeBin, int, 0) // Max time bin for continuous tracking
+AddVariableRTC(dAlpha, float, 0.f)                 // angular size
+AddVariableRTC(assumeConstantBz, signed char, 0)   // Assume a constant magnetic field
+AddVariableRTC(toyMCEventsFlag, signed char, 0)    // events were build with home-made event generator
+AddVariableRTC(continuousTracking, signed char, 0) // Continuous tracking, estimate bz and errors for abs(z) = 125cm during seeding
+AddVariableRTC(dodEdx, signed char, 0)             // Do dEdx computation
+AddVariableRTC(earlyTpcTransform, signed char, 0)  // do Early TPC transformation
+AddVariableRTC(debugLevel, signed char, 0)         // Debug level
 EndConfig()
 
 EndNamespace() // gpu
diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx
@@ -285,7 +285,7 @@ bool GPUChainTracking::ValidateSettings()
     GPUError("OMP Kernels require mergerReadFromTrackerDirectly");
     return false;
   }
-  if (param().par.continuousMaxTimeBin > (int)GPUSettings::TPC_MAX_TF_TIME_BIN) {
+  if (param().continuousMaxTimeBin > (int)GPUSettings::TPC_MAX_TF_TIME_BIN) {
     GPUError("configured max time bin exceeds 256 orbits");
     return false;
   }
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx
@@ -472,7 +472,7 @@ int GPUChainTracking::RunTPCClusterizer_prepare(bool restorePointers)
     mCFContext.reset(new GPUTPCCFChainContext);
   }
   const short maxFragmentLen = GetProcessingSettings().overrideClusterizerFragmentLen;
-  const unsigned int maxAllowedTimebin = param().par.continuousTracking ? std::max<int>(param().par.continuousMaxTimeBin, maxFragmentLen) : TPC_MAX_TIME_BIN_TRIGGERED;
+  const unsigned int maxAllowedTimebin = param().par.continuousTracking ? std::max<int>(param().continuousMaxTimeBin, maxFragmentLen) : TPC_MAX_TIME_BIN_TRIGGERED;
   mCFContext->tpcMaxTimeBin = maxAllowedTimebin;
   const CfFragment fragmentMax{(tpccf::TPCTime)mCFContext->tpcMaxTimeBin + 1, maxFragmentLen};
   mCFContext->prepare(mIOPtrs.tpcZS, fragmentMax);
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx b/GPU/GPUTracking/Global/GPUChainTrackingCompression.cxx
@@ -68,7 +68,7 @@ int GPUChainTracking::RunTPCCompression()
   O->nSliceRows = NSLICES * GPUCA_ROW_COUNT;
   O->nComppressionModes = param().rec.tpc.compressionTypeMask;
   O->solenoidBz = param().bzkG;
-  O->maxTimeBin = param().par.continuousMaxTimeBin;
+  O->maxTimeBin = param().continuousMaxTimeBin;
   size_t outputSize = AllocateRegisteredMemory(Compressor.mMemoryResOutputHost, mSubOutputControls[GPUTrackingOutputs::getIndex(&GPUTrackingOutputs::compressedClusters)]);
   Compressor.mOutputFlat->set(outputSize, *Compressor.mOutput);
   char* hostFlatPtr = (char*)Compressor.mOutput->qTotU; // First array as allocated in GPUTPCCompression::SetPointersCompressedClusters
@@ -253,7 +253,7 @@ int GPUChainTracking::RunTPCDecompression()
     inputGPU.nSliceRows = NSLICES * GPUCA_ROW_COUNT;
     inputGPU.nComppressionModes = param().rec.tpc.compressionTypeMask;
     inputGPU.solenoidBz = param().bzkG;
-    inputGPU.maxTimeBin = param().par.continuousMaxTimeBin;
+    inputGPU.maxTimeBin = param().continuousMaxTimeBin;
     SetupGPUProcessor(&Decompressor, true);
     WriteToConstantMemory(myStep, (char*)&processors()->tpcDecompressor - (char*)processors(), &DecompressorShadow, sizeof(DecompressorShadow), inputStream);
 
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx
@@ -264,7 +264,7 @@ int GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
       }
       GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracks()), outputStream, 0, nullptr, waitEvent);
       waitEvent = nullptr;
-      if (param().par.dodEdx) {
+      if (param().dodEdxDownscaled) {
         GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracksdEdx()), outputStream, 0);
       }
       GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NOutputTrackClusters() * sizeof(*Merger.Clusters()), outputStream, 0);
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx b/GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx
@@ -40,7 +40,7 @@ int GPUChainTracking::RunTRDTracking()
     return 0;
   }
 
-  bool isTriggeredEvent = (param().par.continuousMaxTimeBin == 0);
+  bool isTriggeredEvent = (param().continuousMaxTimeBin == 0);
 
   if (!isTriggeredEvent) {
     Tracker.SetProcessPerTimeFrame(true);
diff --git a/GPU/GPUTracking/Global/GPUChainTrackingTransformation.cxx b/GPU/GPUTracking/Global/GPUChainTrackingTransformation.cxx
@@ -91,7 +91,7 @@ void GPUChainTracking::ConvertNativeToClusterDataLegacy()
   if (tmp != mIOPtrs.clustersNative) {
     *tmp = *mIOPtrs.clustersNative;
   }
-  GPUReconstructionConvert::ConvertNativeToClusterData(mIOMem.clusterNativeAccess.get(), mIOMem.clusterData, mIOPtrs.nClusterData, processors()->calibObjects.fastTransform, param().par.continuousMaxTimeBin);
+  GPUReconstructionConvert::ConvertNativeToClusterData(mIOMem.clusterNativeAccess.get(), mIOMem.clusterData, mIOPtrs.nClusterData, processors()->calibObjects.fastTransform, param().continuousMaxTimeBin);
   for (unsigned int i = 0; i < NSLICES; i++) {
     mIOPtrs.clusterData[i] = mIOMem.clusterData[i].get();
     if (GetProcessingSettings().registerStandaloneInputMemory) {
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
@@ -325,7 +325,7 @@ void* GPUTPCGMMerger::SetPointersRefitScratch2(void* mem)
 void* GPUTPCGMMerger::SetPointersOutput(void* mem)
 {
   computePointerWithAlignment(mem, mOutputTracks, mNMaxTracks);
-  if (mRec->GetParam().par.dodEdx) {
+  if (mRec->GetParam().dodEdxDownscaled) {
     computePointerWithAlignment(mem, mOutputTracksdEdx, mNMaxTracks);
   }
   computePointerWithAlignment(mem, mClusters, mNMaxOutputTrackClusters);
@@ -463,7 +463,7 @@ GPUd() int GPUTPCGMMerger::RefitSliceTrack(GPUTPCGMSliceTrack& sliceTrack, const
   trk.SinPhi() = inTrack->Param().GetSinPhi();
   trk.DzDs() = inTrack->Param().GetDzDs();
   trk.QPt() = inTrack->Param().GetQPt();
-  trk.TZOffset() = Param().par.earlyTpcTransform ? inTrack->Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(slice, inTrack->Param().GetZOffset(), Param().par.continuousMaxTimeBin);
+  trk.TZOffset() = Param().par.earlyTpcTransform ? inTrack->Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(slice, inTrack->Param().GetZOffset(), Param().continuousMaxTimeBin);
   trk.ShiftZ(this, slice, sliceTrack.ClusterZT0(), sliceTrack.ClusterZTN(), inTrack->Param().GetX(), inTrack->Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
   sliceTrack.SetX2(0.f);
   for (int way = 0; way < 2; way++) {
@@ -2125,7 +2125,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersInit(int nBlocks, int nThreads, int iBlo
     const float qptabs = CAMath::Abs(p.GetQPt());
     if (trk.NClusters() && qptabs * Param().qptB5Scaler > 5.f && qptabs * Param().qptB5Scaler <= lowPtThresh) {
       const int slice = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].slice;
-      const float refz = p.GetZ() + (Param().par.earlyTpcTransform ? p.GetTZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(slice, p.GetTZOffset(), Param().par.continuousMaxTimeBin)) + (trk.CSide() ? -100 : 100);
+      const float refz = p.GetZ() + (Param().par.earlyTpcTransform ? p.GetTZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(slice, p.GetTZOffset(), Param().continuousMaxTimeBin)) + (trk.CSide() ? -100 : 100);
       float sinA, cosA;
       CAMath::SinCos(trk.GetAlpha(), sinA, cosA);
       float gx = cosA * p.GetX() - sinA * p.GetY();
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx
@@ -44,7 +44,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::prepare>(int nBlocks, i
 
   constexpr unsigned char flagsReject = getFlagsReject();
   const unsigned int flagsRequired = getFlagsRequired(merger.Param().rec);
-  bool cutOnTrackdEdx = merger.Param().par.dodEdx && merger.Param().rec.tpc.minTrackdEdxMax2Tot > 0.f;
+  bool cutOnTrackdEdx = merger.Param().par.dodEdx && merger.Param().dodEdxDownscaled && merger.Param().rec.tpc.minTrackdEdxMax2Tot > 0.f;
 
   GPUTPCGMMerger::tmpSort* GPUrestrict() trackSort = merger.TrackSortO2();
   uint2* GPUrestrict() tmpData = merger.ClusRefTmp();
@@ -146,7 +146,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int nBlocks, in
 
     oTrack.setChi2(tracks[i].GetParam().GetChi2());
     auto& outerPar = tracks[i].OuterParam();
-    if (merger.Param().par.dodEdx) {
+    if (merger.Param().par.dodEdx && merger.Param().dodEdxDownscaled) {
       oTrack.setdEdx(tracksdEdx[i]);
     }
 
@@ -163,7 +163,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int nBlocks, in
        outerPar.C[6], outerPar.C[7], outerPar.C[8], outerPar.C[9], outerPar.C[10], outerPar.C[11],
        outerPar.C[12], outerPar.C[13], outerPar.C[14]}));
 
-    if (merger.Param().par.dodEdx && merger.Param().rec.tpc.enablePID) {
+    if (merger.Param().par.dodEdx && merger.Param().dodEdxDownscaled && merger.Param().rec.tpc.enablePID) {
       PIDResponse pidResponse{};
       auto pid = pidResponse.getMostProbablePID(oTrack, merger.Param().rec.tpc.PID_EKrangeMin, merger.Param().rec.tpc.PID_EKrangeMax, merger.Param().rec.tpc.PID_EPrangeMin, merger.Param().rec.tpc.PID_EPrangeMax, merger.Param().rec.tpc.PID_EDrangeMin, merger.Param().rec.tpc.PID_EDrangeMax, merger.Param().rec.tpc.PID_ETrangeMin, merger.Param().rec.tpc.PID_ETrangeMax, merger.Param().rec.tpc.PID_useNsigma, merger.Param().rec.tpc.PID_sigma);
       auto pidRemap = merger.Param().rec.tpc.PID_remap[pid];
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSliceTrack.cxx b/GPU/GPUTracking/Merger/GPUTPCGMSliceTrack.cxx
@@ -40,7 +40,7 @@ GPUd() void GPUTPCGMSliceTrack::Set(const GPUTPCGMMerger* merger, const GPUTPCTr
   if (merger->Param().par.earlyTpcTransform) {
     mTZOffset = t.GetZOffset();
   } else {
-    mTZOffset = merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(slice, t.GetZOffset(), merger->Param().par.continuousMaxTimeBin);
+    mTZOffset = merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(slice, t.GetZOffset(), merger->Param().continuousMaxTimeBin);
   }
   mNClusters = sliceTr->NHits();
 }
@@ -329,7 +329,7 @@ GPUd() bool GPUTPCGMSliceTrack::TransportToX(GPUTPCGMMerger* merger, float x, fl
   if (merger->Param().par.earlyTpcTransform) {
     b.SetZOffsetLinear(mTZOffset);
   } else {
-    b.SetZOffsetLinear(merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(mSlice, mTZOffset, merger->Param().par.continuousMaxTimeBin));
+    b.SetZOffsetLinear(merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(mSlice, mTZOffset, merger->Param().continuousMaxTimeBin));
   }
 
   if (!doCov) {
@@ -485,7 +485,7 @@ GPUd() bool GPUTPCGMSliceTrack::TransportToXAlpha(GPUTPCGMMerger* merger, float
   if (merger->Param().par.earlyTpcTransform) {
     b.SetZOffsetLinear(mTZOffset);
   } else {
-    b.SetZOffsetLinear(merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(mSlice, mTZOffset, merger->Param().par.continuousMaxTimeBin));
+    b.SetZOffsetLinear(merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(mSlice, mTZOffset, merger->Param().continuousMaxTimeBin));
   }
 
   b.SetCov(0, c00 + h2 * h2c22 + h4 * h4c44 + 2.f * (h2 * c20ph4c42 + h4 * c40));
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx
diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx
diff --git a/GPU/GPUTracking/Standalone/tools/GPUExtractPbPbCollision.h b/GPU/GPUTracking/Standalone/tools/GPUExtractPbPbCollision.h
diff --git a/GPU/GPUTracking/TPCConvert/GPUTPCConvertImpl.h b/GPU/GPUTracking/TPCConvert/GPUTPCConvertImpl.h
diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx
diff --git a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx
diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx

Original file line number	Diff line number	Diff line change
`@@ -119,7 +119,7 @@ void GPUParam::SetDefaults(float solenoidBz)`
`119`	`119`	`par.assumeConstantBz = false;`
`120`	`120`	`par.toyMCEventsFlag = false;`
`121`	`121`	`par.continuousTracking = false;`
`122`		`- par.continuousMaxTimeBin = 0;`
	`122`	`+ continuousMaxTimeBin = 0;`
`123`	`123`	`par.debugLevel = 0;`
`124`	`124`	`par.earlyTpcTransform = false;`
`125`	`125`	`}`
`@@ -131,7 +131,7 @@ void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessi`
`131`	`131`	`par.assumeConstantBz = g->constBz;`
`132`	`132`	`par.toyMCEventsFlag = g->homemadeEvents;`
`133`	`133`	`par.continuousTracking = g->continuousMaxTimeBin != 0;`
`134`		`- par.continuousMaxTimeBin = g->continuousMaxTimeBin == -1 ? GPUSettings::TPC_MAX_TF_TIME_BIN : g->continuousMaxTimeBin;`
	`134`	`+ continuousMaxTimeBin = g->continuousMaxTimeBin == -1 ? GPUSettings::TPC_MAX_TF_TIME_BIN : g->continuousMaxTimeBin;`
`135`	`135`	`}`
`136`	`136`	`par.earlyTpcTransform = rec.tpc.forceEarlyTransform == -1 ? (!par.continuousTracking) : rec.tpc.forceEarlyTransform;`
`137`	`137`	`qptB5Scaler = CAMath::Abs(bzkG) > 0.1f ? CAMath::Abs(bzkG) / 5.006680f : 1.f; // Repeat here, since passing in g is optional`
`@@ -140,9 +140,9 @@ void GPUParam::UpdateSettings(const GPUSettingsGRP* g, const GPUSettingsProcessi`
`140`	`140`	`UpdateRun3ClusterErrors(p->param.tpcErrorParamY, p->param.tpcErrorParamZ);`
`141`	`141`	`}`
`142`	`142`	`if (w) {`
`143`		`- par.dodEdx = w->steps.isSet(GPUDataTypes::RecoStep::TPCdEdx);`
	`143`	`+ par.dodEdx = dodEdxDownscaled = w->steps.isSet(GPUDataTypes::RecoStep::TPCdEdx);`
`144`	`144`	`if (par.dodEdx && p && p->tpcDownscaledEdx != 0) {`
`145`		`- par.dodEdx = (rand() % 100) < p->tpcDownscaledEdx;`
	`145`	`+ dodEdxDownscaled = (rand() % 100) < p->tpcDownscaledEdx;`
`146`	`146`	`}`
`147`	`147`	`}`
`148`	`148`	`}`
Original file line number	Diff line number	Diff line change
`@@ -28,7 +28,7 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste`
`28`	`28`	`CompressedClusters& GPUrestrict() cmprClusters = decompressor.mInputGPU;`
`29`	`29`	`const GPUParam& GPUrestrict() param = processors.param;`
`30`	`30`
`31`		`- const unsigned int maxTime = (param.par.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1;`
	`31`	`+ const unsigned int maxTime = (param.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1;`
`32`	`32`
`33`	`33`	`for (int i = trackStart + get_global_id(0); i < trackEnd; i += get_global_size(0)) {`
`34`	`34`	`decompressTrack(cmprClusters, param, maxTime, i, decompressor.mAttachedClustersOffsets[i], decompressor);`
`@@ -85,7 +85,7 @@ GPUdii() void GPUTPCDecompressionKernels::decompressTrack(CompressedClusters& cm`
`85`	`85`	`pad = param.tpcGeometry.NPads(row) * ClusterNative::scalePadPacked - 1;`
`86`	`86`	`}`
`87`	`87`	`}`
`88`		`- if (param.par.continuousMaxTimeBin > 0 && time >= maxTime) {`
	`88`	`+ if (param.continuousMaxTimeBin > 0 && time >= maxTime) {`
`89`	`89`	`if (time >= 0xFFFFFF - 544768) { // Constant 544768 = (2^23 - LHCMAXBUNCHES(3564) * MAXORBITS(256) * scaleTimePacked(64) / BCPERTIMEBIN(8)) / 2)`
`90`	`90`	`time = 0;`
`91`	`91`	`} else {`
`@@ -152,8 +152,8 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste`
`152`	`152`	`if (t < 0) {`
`153`	`153`	`t = 0;`
`154`	`154`	`}`
`155`		`- if (processors.param.par.continuousMaxTimeBin > 0 && t > processors.param.par.continuousMaxTimeBin) {`
`156`		`- t = processors.param.par.continuousMaxTimeBin;`
	`155`	`+ if (processors.param.continuousMaxTimeBin > 0 && t > processors.param.continuousMaxTimeBin) {`
	`156`	`+ t = processors.param.continuousMaxTimeBin;`
`157`	`157`	`}`
`158`	`158`	`cl.setTime(t);`
`159`	`159`	`}`
Original file line number	Diff line number	Diff line change
`@@ -43,7 +43,7 @@ int TPCClusterDecompressor::decompress(const CompressedClusters* clustersCompres`
`43`	`43`	`if (clustersCompressed->nTracks && clustersCompressed->solenoidBz != -1e6f && clustersCompressed->solenoidBz != param.bzkG) {`
`44`	`44`	`throw std::runtime_error("Configured solenoid Bz does not match value used for track model encoding");`
`45`	`45`	`}`
`46`		`- if (clustersCompressed->nTracks && clustersCompressed->maxTimeBin != -1e6 && clustersCompressed->maxTimeBin != param.par.continuousMaxTimeBin) {`
	`46`	`+ if (clustersCompressed->nTracks && clustersCompressed->maxTimeBin != -1e6 && clustersCompressed->maxTimeBin != param.continuousMaxTimeBin) {`
`47`	`47`	`throw std::runtime_error("Configured max time bin does not match value used for track model encoding");`
`48`	`48`	`}`
`49`	`49`	`std::vector<ClusterNative> clusters[NSLICES][GPUCA_ROW_COUNT];`
`@@ -52,7 +52,7 @@ int TPCClusterDecompressor::decompress(const CompressedClusters* clustersCompres`
`52`	`52`	`(&locks[0][0])[i].clear();`
`53`	`53`	`}`
`54`	`54`	`unsigned int offset = 0, lasti = 0;`
`55`		`- const unsigned int maxTime = param.par.continuousMaxTimeBin > 0 ? ((param.par.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1) : TPC_MAX_TIME_BIN_TRIGGERED;`
	`55`	`+ const unsigned int maxTime = param.continuousMaxTimeBin > 0 ? ((param.continuousMaxTimeBin + 1) * ClusterNative::scaleTimePacked - 1) : TPC_MAX_TIME_BIN_TRIGGERED;`
`56`	`56`	`GPUCA_OPENMP(parallel for firstprivate(offset, lasti))`
`57`	`57`	`for (unsigned int i = 0; i < clustersCompressed->nTracks; i++) {`
`58`	`58`	`if (i < lasti) {`
`@@ -99,8 +99,8 @@ int TPCClusterDecompressor::decompress(const CompressedClusters* clustersCompres`
`99`	`99`	`if (t < 0) {`
`100`	`100`	`t = 0;`
`101`	`101`	`}`
`102`		`- if (param.par.continuousMaxTimeBin > 0 && t > param.par.continuousMaxTimeBin) {`
`103`		`- t = param.par.continuousMaxTimeBin;`
	`102`	`+ if (param.continuousMaxTimeBin > 0 && t > param.continuousMaxTimeBin) {`
	`103`	`+ t = param.continuousMaxTimeBin;`
`104`	`104`	`}`
`105`	`105`	`cl.setTime(t);`
`106`	`106`	`}`
Original file line number	Diff line number	Diff line change
`@@ -101,7 +101,7 @@ inline void TPCClusterDecompressor::decompressTrack(const CompressedClusters* cl`
`101`	`101`	`pad = param.tpcGeometry.NPads(row) * ClusterNative::scalePadPacked - 1;`
`102`	`102`	`}`
`103`	`103`	`}`
`104`		`- if (param.par.continuousMaxTimeBin > 0 && time >= maxTime) {`
	`104`	`+ if (param.continuousMaxTimeBin > 0 && time >= maxTime) {`
`105`	`105`	`if (time >= 0xFFFFFF - 544768) { // Constant 544768 = (2^23 - LHCMAXBUNCHES(3564) * MAXORBITS(256) * scaleTimePacked(64) / BCPERTIMEBIN(8)) / 2)`
`106`	`106`	`time = 0;`
`107`	`107`	`} else {`
Original file line number	Diff line number	Diff line change
`@@ -22,7 +22,7 @@ GPUd() unsigned int GPUTPCClusterOccupancyMapBin::getNBins(const GPUParam& param`
`22`	`22`	`if (param.rec.tpc.occupancyMapTimeBins == 0) {`
`23`	`23`	`return 0;`
`24`	`24`	`}`
`25`		`- unsigned int maxTimeBin = param.par.continuousTracking ? param.par.continuousMaxTimeBin : TPC_MAX_TIME_BIN_TRIGGERED;`
	`25`	`+ unsigned int maxTimeBin = param.par.continuousTracking ? param.continuousMaxTimeBin : TPC_MAX_TIME_BIN_TRIGGERED;`
`26`	`26`	`return (maxTimeBin + param.rec.tpc.occupancyMapTimeBins) / param.rec.tpc.occupancyMapTimeBins; // Not -1, since maxTimeBin is allowed`
`27`	`27`	`}`
`28`	`28`
Original file line number	Diff line number	Diff line change
`@@ -285,7 +285,7 @@ bool GPUChainTracking::ValidateSettings()`
`285`	`285`	`GPUError("OMP Kernels require mergerReadFromTrackerDirectly");`
`286`	`286`	`return false;`
`287`	`287`	`}`
`288`		`- if (param().par.continuousMaxTimeBin > (int)GPUSettings::TPC_MAX_TF_TIME_BIN) {`
	`288`	`+ if (param().continuousMaxTimeBin > (int)GPUSettings::TPC_MAX_TF_TIME_BIN) {`
`289`	`289`	`GPUError("configured max time bin exceeds 256 orbits");`
`290`	`290`	`return false;`
`291`	`291`	`}`
Original file line number	Diff line number	Diff line change
`@@ -472,7 +472,7 @@ int GPUChainTracking::RunTPCClusterizer_prepare(bool restorePointers)`
`472`	`472`	`mCFContext.reset(new GPUTPCCFChainContext);`
`473`	`473`	`}`
`474`	`474`	`const short maxFragmentLen = GetProcessingSettings().overrideClusterizerFragmentLen;`
`475`		`- const unsigned int maxAllowedTimebin = param().par.continuousTracking ? std::max<int>(param().par.continuousMaxTimeBin, maxFragmentLen) : TPC_MAX_TIME_BIN_TRIGGERED;`
	`475`	`+ const unsigned int maxAllowedTimebin = param().par.continuousTracking ? std::max<int>(param().continuousMaxTimeBin, maxFragmentLen) : TPC_MAX_TIME_BIN_TRIGGERED;`
`476`	`476`	`mCFContext->tpcMaxTimeBin = maxAllowedTimebin;`
`477`	`477`	`const CfFragment fragmentMax{(tpccf::TPCTime)mCFContext->tpcMaxTimeBin + 1, maxFragmentLen};`
`478`	`478`	`mCFContext->prepare(mIOPtrs.tpcZS, fragmentMax);`