diff --git a/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx b/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx index fc707f433c8c5..35b76715cbc28 100644 --- 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>("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 index d9518b903d98f..a0d2f9c297f7c 100644 --- 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 index 852f52bf68862..ef1cccf58cdbf 100644 --- 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 index 7a58cad637514..be3abfbdfa598 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx @@ -28,7 +28,7 @@ GPUdii() void GPUTPCDecompressionKernels::Thread 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 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 index 61e3392af0f03..d829dc978c5ef 100644 --- 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 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 index f83b539d80cfa..f0cd648e0f49e 100644 --- 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 index 96c288d86b23d..9b388bb0260e5 100644 --- 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 index 610cfe1627b2f..26a6168faecbb 100644 --- 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 index fb5c6974e03c3..ba17b3e3ead54 100644 --- 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 index c26968513611a..bcf4310a333f6 100644 --- 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(param().par.continuousMaxTimeBin, maxFragmentLen) : TPC_MAX_TIME_BIN_TRIGGERED; + const unsigned int maxAllowedTimebin = param().par.continuousTracking ? std::max(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 index 3c06bb7a3a3c5..83bcb20ff6e8b 100644 --- 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 index 185de2f804acf..cd6de46b95c7a 100644 --- 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 index 39b9492417400..3d16b6693cade 100644 --- 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 index 7c8b77d9fe36c..84aef65146ca3 100644 --- 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 index 671ec9a7794a0..379be50b09caf 100644 --- 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 index fe3157bfae3ae..de9b0aa0a2038 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx @@ -44,7 +44,7 @@ GPUdii() void GPUTPCGMO2Output::Thread(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(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(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 index b6f2c72a209e0..c452d2f961979 100644 --- 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 index 249e99f240ca0..5e1a6a37cb13a 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx @@ -212,7 +212,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int iT continue; } } else if (allowModification && lastRow != 255 && CAMath::Abs(cluster.row - lastRow) > 1) { - bool dodEdx = merger->Param().par.dodEdx && merger->Param().rec.tpc.adddEdxSubThresholdClusters && iWay == nWays - 1 && CAMath::Abs(cluster.row - lastRow) == 2 && cluster.leg == clusters[maxN - 1].leg; + bool dodEdx = merger->Param().par.dodEdx && merger->Param().dodEdxDownscaled && merger->Param().rec.tpc.adddEdxSubThresholdClusters && iWay == nWays - 1 && CAMath::Abs(cluster.row - lastRow) == 2 && cluster.leg == clusters[maxN - 1].leg; dodEdx = AttachClustersPropagate(merger, cluster.slice, lastRow, cluster.row, iTrk, cluster.leg == clusters[maxN - 1].leg, prop, inFlyDirection, GPUCA_MAX_SIN_PHI, dodEdx); if (dodEdx) { dEdx.fillSubThreshold(lastRow - 1, param); @@ -362,7 +362,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int iT CADEBUG(printf("Reinit linearization\n")); prop.SetTrack(this, prop.GetAlpha()); } - if (merger->Param().par.dodEdx && iWay == nWays - 1 && cluster.leg == clusters[maxN - 1].leg && !(clusterState & GPUTPCGMMergedTrackHit::flagEdge)) { + if (merger->Param().par.dodEdx && merger->Param().dodEdxDownscaled && iWay == nWays - 1 && cluster.leg == clusters[maxN - 1].leg && !(clusterState & GPUTPCGMMergedTrackHit::flagEdge)) { float qtot = 0, qmax = 0, pad = 0, relTime = 0; const int clusterCount = (ihit - ihitMergeFirst) * wayDirection + 1; for (int iTmp = ihitMergeFirst; iTmp != ihit + wayDirection; iTmp += wayDirection) { @@ -413,7 +413,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int iT // TODO: we have looping tracks here with 0 accepted clusters in the primary leg. In that case we should refit the track using only the primary leg. - if (merger->Param().par.dodEdx) { + if (merger->Param().par.dodEdx && merger->Param().dodEdxDownscaled) { dEdx.computedEdx(merger->OutputTracksdEdx()[iTrk], param); } Alpha = prop.GetAlpha(); @@ -574,7 +574,7 @@ GPUd() float GPUTPCGMTrackParam::AttachClusters(const GPUTPCGMMerger* GPUrestric return -1e6f; } - const float zOffset = Merger->Param().par.earlyTpcTransform ? ((Merger->OutputTracks()[iTrack].CSide() ^ (slice >= 18)) ? -mTZOffset : mTZOffset) : Merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(slice, mTZOffset, Merger->Param().par.continuousMaxTimeBin); + const float zOffset = Merger->Param().par.earlyTpcTransform ? ((Merger->OutputTracks()[iTrack].CSide() ^ (slice >= 18)) ? -mTZOffset : mTZOffset) : Merger->GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(slice, mTZOffset, Merger->Param().continuousMaxTimeBin); const float y0 = row.Grid().YMin(); const float stepY = row.HstepY(); const float z0 = row.Grid().ZMin() - zOffset; // We can use our own ZOffset, since this is only used temporarily anyway diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx index 85275e19c844b..31a2c5828fe7d 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx +++ b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx @@ -112,7 +112,7 @@ void* GPUTPCSliceData::SetPointersRows(void* mem) GPUd() void GPUTPCSliceData::GetMaxNBins(GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * mem, GPUTPCRow* GPUrestrict() row, int& maxY, int& maxZ) { maxY = row->mMaxY * 2.f / GPUCA_MIN_BIN_SIZE + 1; - maxZ = (mem->param.par.continuousMaxTimeBin > 0 ? (mem->calibObjects.fastTransformHelper->getCorrMap()->convTimeToZinTimeFrame(0, 0, mem->param.par.continuousMaxTimeBin)) : mem->param.tpcGeometry.TPCLength()) + 50; + maxZ = (mem->param.continuousMaxTimeBin > 0 ? (mem->calibObjects.fastTransformHelper->getCorrMap()->convTimeToZinTimeFrame(0, 0, mem->param.continuousMaxTimeBin)) : mem->param.tpcGeometry.TPCLength()) + 50; maxZ = maxZ / GPUCA_MIN_BIN_SIZE + 1; } diff --git a/GPU/GPUTracking/Standalone/tools/GPUExtractPbPbCollision.h b/GPU/GPUTracking/Standalone/tools/GPUExtractPbPbCollision.h index f1946496eef55..e0cd756d32651 100644 --- a/GPU/GPUTracking/Standalone/tools/GPUExtractPbPbCollision.h +++ b/GPU/GPUTracking/Standalone/tools/GPUExtractPbPbCollision.h @@ -14,18 +14,18 @@ static void GPUExtractPbPbCollision(GPUParam& param, GPUTrackingInOutPointers& ioPtrs) { - std::vector counts(param.par.continuousMaxTimeBin + 1); - std::vector sums(param.par.continuousMaxTimeBin + 1); - std::vector countsTracks(param.par.continuousMaxTimeBin + 1); - std::vector sumsTracks(param.par.continuousMaxTimeBin + 1); - std::vector mask(param.par.continuousMaxTimeBin + 1); + std::vector counts(param.continuousMaxTimeBin + 1); + std::vector sums(param.continuousMaxTimeBin + 1); + std::vector countsTracks(param.continuousMaxTimeBin + 1); + std::vector sumsTracks(param.continuousMaxTimeBin + 1); + std::vector mask(param.continuousMaxTimeBin + 1); const int driftlength = 520; const bool checkAfterGlow = true; const int afterGlowLength = checkAfterGlow ? 8000 : 0; for (unsigned int i = 0; i < ioPtrs.clustersNative->nClustersTotal; i++) { int time = ioPtrs.clustersNative->clustersLinear[i].getTime(); - if (time < 0 || time > param.par.continuousMaxTimeBin) { - fprintf(stderr, "Invalid time %d > %d\n", time, param.par.continuousMaxTimeBin); + if (time < 0 || time > param.continuousMaxTimeBin) { + fprintf(stderr, "Invalid time %d > %d\n", time, param.continuousMaxTimeBin); throw std::runtime_error("Invalid Time"); } counts[time]++; @@ -35,19 +35,19 @@ static void GPUExtractPbPbCollision(GPUParam& param, GPUTrackingInOutPointers& i continue; } int time = ioPtrs.mergedTracks[i].GetParam().GetTZOffset(); - if (time < 0 || time > param.par.continuousMaxTimeBin) { + if (time < 0 || time > param.continuousMaxTimeBin) { continue; } countsTracks[time]++; } int first = 0, last = 0; - for (int i = driftlength; i < param.par.continuousMaxTimeBin; i++) { + for (int i = driftlength; i < param.continuousMaxTimeBin; i++) { if (counts[i]) { first = i; break; } } - for (int i = param.par.continuousMaxTimeBin + 1 - driftlength; i > 0; i--) { + for (int i = param.continuousMaxTimeBin + 1 - driftlength; i > 0; i--) { if (counts[i - 1]) { last = i; break; diff --git a/GPU/GPUTracking/TPCConvert/GPUTPCConvertImpl.h b/GPU/GPUTracking/TPCConvert/GPUTPCConvertImpl.h index b791a0aa08abe..e8d777a3b23c1 100644 --- a/GPU/GPUTracking/TPCConvert/GPUTPCConvertImpl.h +++ b/GPU/GPUTracking/TPCConvert/GPUTPCConvertImpl.h @@ -31,7 +31,7 @@ class GPUTPCConvertImpl GPUd() static void convert(const GPUConstantMem& GPUrestrict() cm, int slice, int row, float pad, float time, float& GPUrestrict() x, float& GPUrestrict() y, float& GPUrestrict() z) { if (cm.param.par.continuousTracking) { - cm.calibObjects.fastTransformHelper->getCorrMap()->TransformInTimeFrame(slice, row, pad, time, x, y, z, cm.param.par.continuousMaxTimeBin); + cm.calibObjects.fastTransformHelper->getCorrMap()->TransformInTimeFrame(slice, row, pad, time, x, y, z, cm.param.continuousMaxTimeBin); } else { cm.calibObjects.fastTransformHelper->Transform(slice, row, pad, time, x, y, z); } @@ -39,7 +39,7 @@ class GPUTPCConvertImpl GPUd() static void convert(const TPCFastTransform& GPUrestrict() transform, const GPUParam& GPUrestrict() param, int slice, int row, float pad, float time, float& GPUrestrict() x, float& GPUrestrict() y, float& GPUrestrict() z) { if (param.par.continuousTracking) { - transform.TransformInTimeFrame(slice, row, pad, time, x, y, z, param.par.continuousMaxTimeBin); + transform.TransformInTimeFrame(slice, row, pad, time, x, y, z, param.continuousMaxTimeBin); } else { transform.Transform(slice, row, pad, time, x, y, z); } diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 8adfb8441eee2..0b0eed697de2d 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -518,12 +518,12 @@ void GPUDisplay::DrawFinal(int iSlice, int /*iCol*/, GPUTPCGMPropagator* prop, s auto cl = mIOPtrs->mergedTrackHits[track->FirstClusterRef() + lastCluster]; const auto& cln = mIOPtrs->clustersNative->clustersLinear[cl.num]; GPUTPCConvertImpl::convert(*mCalib->fastTransform, *mParam, cl.slice, cl.row, cln.getPad(), cln.getTime(), x, y, z); - ZOffset = mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(iSlice, track->GetParam().GetTZOffset(), mParam->par.continuousMaxTimeBin); + ZOffset = mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(iSlice, track->GetParam().GetTZOffset(), mParam->continuousMaxTimeBin); } else { uint8_t sector, row; auto cln = track->getCluster(mIOPtrs->outputClusRefsTPCO2, lastCluster, *mIOPtrs->clustersNative, sector, row); GPUTPCConvertImpl::convert(*mCalib->fastTransform, *mParam, sector, row, cln.getPad(), cln.getTime(), x, y, z); - ZOffset = mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(sector, track->getTime0(), mParam->par.continuousMaxTimeBin); + ZOffset = mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(sector, track->getTime0(), mParam->continuousMaxTimeBin); } } } else { @@ -553,7 +553,7 @@ void GPUDisplay::DrawFinal(int iSlice, int /*iCol*/, GPUTPCGMPropagator* prop, s #ifdef GPUCA_TPC_GEOMETRY_O2 trkParam.Set(mclocal[0], mclocal[1], mc.z, mclocal[2], mclocal[3], mc.pZ, -charge); // TODO: DR: unclear to me why we need -charge here if (mParam->par.continuousTracking) { - ZOffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, mc.t0, mParam->par.continuousMaxTimeBin)) * (mc.z < 0 ? -1 : 1); + ZOffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, mc.t0, mParam->continuousMaxTimeBin)) * (mc.z < 0 ? -1 : 1); } #else if (fabsf(mc.z) > GPUTPCGeometry::TPCLength()) { diff --git a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx index 2b7baec173b61..b1e6dfe21fb97 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx @@ -184,7 +184,7 @@ void GPUDisplay::DrawGLScene_updateEventData() trdTriggerRecord++; #ifdef GPUCA_HAVE_O2HEADERS float trdTime = mIOPtrs->trdTriggerTimes[trdTriggerRecord] * 1e3 / o2::constants::lhc::LHCBunchSpacingNS / o2::tpc::constants::LHCBCPERTIMEBIN; - trdZoffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, trdTime, mParam->par.continuousMaxTimeBin)); + trdZoffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, trdTime, mParam->continuousMaxTimeBin)); #endif } const auto& sp = mIOPtrs->trdSpacePoints[i]; @@ -219,7 +219,7 @@ void GPUDisplay::DrawGLScene_updateEventData() float ZOffset = 0; if (mParam->par.continuousTracking) { float tofTime = mIOPtrs->tofClusters[i].getTime() * 1e-3 / o2::constants::lhc::LHCBunchSpacingNS / o2::tpc::constants::LHCBCPERTIMEBIN; - ZOffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, tofTime, mParam->par.continuousMaxTimeBin)); + ZOffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, tofTime, mParam->continuousMaxTimeBin)); ptr->z += ptr->z > 0 ? ZOffset : -ZOffset; } if (fabsf(ptr->z) > mMaxClusterZ) { @@ -247,7 +247,7 @@ void GPUDisplay::DrawGLScene_updateEventData() if (mParam->par.continuousTracking) { o2::InteractionRecord startIR = o2::InteractionRecord(0, mIOPtrs->settingsTF && mIOPtrs->settingsTF->hasTfStartOrbit ? mIOPtrs->settingsTF->tfStartOrbit : 0); float itsROFtime = mIOPtrs->itsClusterROF[j].getBCData().differenceInBC(startIR) / (float)o2::tpc::constants::LHCBCPERTIMEBIN; - ZOffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, itsROFtime + itsROFhalfLen, mParam->par.continuousMaxTimeBin)); + ZOffset = fabsf(mCalib->fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(0, itsROFtime + itsROFhalfLen, mParam->continuousMaxTimeBin)); } if (i != mIOPtrs->itsClusterROF[j].getFirstEntry()) { throw std::runtime_error("Inconsistent ITS data, number of clusters does not match ROF content"); diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx index 6b505e133f88e..85b0b49305fc9 100644 --- a/GPU/GPUTracking/qa/GPUQA.cxx +++ b/GPU/GPUTracking/qa/GPUQA.cxx @@ -1363,7 +1363,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx if (tracksExternal) { return param.GetZ(); } - if (!mParam->par.continuousMaxTimeBin) { + if (!mParam->continuousMaxTimeBin) { return param.GetZ() - mc1.z; } #ifdef GPUCA_TPC_GEOMETRY_O2