Skip to content

Commit 64d520e

Browse files
davidrohrdavidrohr
committed
GPU TPC: Replace rejectQPt and maxTrackQPt by B=0.5T reference values rejectQPtB5 and maxTrackQPtB5
1 parent 09c2bf8 commit 64d520e

8 files changed

Lines changed: 25 additions & 30 deletions

File tree

‎GPU/GPUTracking/Base/GPUParam.cxx‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -139,10 +139,6 @@ void GPUParam::UpdateGRPSettings(const GPUSettingsGRP* g, const GPUSettingsProce
139139
if (p) {
140140
par.debugLevel = p->debugLevel;
141141
par.resetTimers = p->resetTimers;
142-
if (p->automaticQPtThresholds) {
143-
rec.maxTrackQPt = 1.f / GPUCA_MIN_TRACK_PTB5_DEFAULT / par.qptB5Scaler;
144-
rec.tpc.rejectQPt = 1.f / GPUCA_MIN_TRACK_PTB5_REJECT / par.qptB5Scaler;
145-
}
146142
}
147143
}
148144

‎GPU/GPUTracking/Benchmark/standalone.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -456,7 +456,7 @@ int SetupReconstruction()
456456
procSet.runCompressionStatistics = 0;
457457
recSet.tpc.disableRefitAttachment = 0xFF;
458458
recSet.tpc.loopInterpolationInExtraPass = 0;
459-
recSet.maxTrackQPt = CAMath::Min(recSet.maxTrackQPt, recSet.tpc.rejectQPt);
459+
recSet.maxTrackQPtB5 = CAMath::Min(recSet.maxTrackQPtB5, recSet.tpc.rejectQPtB5);
460460
recSet.useMatLUT = true;
461461
recAsync->SetSettings(&grp, &recSet, &procSet, &steps);
462462
}

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
4040
if (!trk.OK()) {
4141
continue;
4242
}
43-
bool rejectTrk = CAMath::Abs(trk.GetParam().GetQPt()) > processors.param.rec.tpc.rejectQPt || trk.MergedLooper();
43+
bool rejectTrk = CAMath::Abs(trk.GetParam().GetQPt() * processors.param.par.qptB5Scaler) > processors.param.rec.tpc.rejectQPtB5 || trk.MergedLooper();
4444
unsigned int nClustersStored = 0;
4545
CompressedClustersPtrs& GPUrestrict() c = compressor.mPtrs;
4646
unsigned int lastRow = 0, lastSlice = 0; // BUG: These should be unsigned char, but then CUDA breaks
@@ -222,7 +222,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
222222
}
223223
int id = attach & gputpcgmmergertypes::attachTrackMask;
224224
auto& trk = ioPtrs.mergedTracks[id];
225-
if (CAMath::Abs(trk.GetParam().GetQPt()) > processors.param.rec.tpc.rejectQPt || trk.MergedLooper()) {
225+
if (CAMath::Abs(trk.GetParam().GetQPt() * processors.param.par.qptB5Scaler) > processors.param.rec.tpc.rejectQPtB5 || trk.MergedLooper()) {
226226
break;
227227
}
228228
}

‎GPU/GPUTracking/Definitions/GPUSettingsList.h‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@ BeginNamespace(gpu)
3636
// Settings concerning the reconstruction
3737
// There must be no bool in here, use char, as sizeof(bool) is compiler dependent and fails on GPUs!!!!!!
3838
BeginSubConfig(GPUSettingsRecTPC, tpc, configStandalone.rec, "RECTPC", 0, "Reconstruction settings", rec_tpc)
39-
AddOptionRTC(rejectQPt, float, 1.f / GPUCA_MIN_TRACK_PTB5_REJECT, "", 0, "QPt threshold to reject clusters of TPC tracks (Inverse Pt!!!)")
39+
AddOptionRTC(rejectQPtB5, float, 1.f / GPUCA_MIN_TRACK_PTB5_REJECT, "", 0, "QPt threshold to reject clusters of TPC tracks (Inverse Pt, scaled to B=0.5T!!!)")
4040
AddOptionRTC(hitPickUpFactor, float, 2., "", 0, "multiplier for the chi2 window for hit pick up procedure")
4141
AddOptionRTC(neighboursSearchArea, float, 3., "", 0, "area in cm for the search of neighbours")
4242
AddOptionRTC(clusterError2CorrectionY, float, 1., "", 0, "correction for the squared cluster error during tracking")
@@ -96,15 +96,15 @@ AddHelp("help", 'h')
9696
EndConfig()
9797

9898
BeginSubConfig(GPUSettingsRec, rec, configStandalone, "REC", 0, "Reconstruction settings", rec)
99-
AddOptionRTC(maxTrackQPt, float, 1.f / GPUCA_MIN_TRACK_PTB5_DEFAULT, "", 0, "required max Q/Pt (==min Pt) of tracks")
99+
AddOptionRTC(maxTrackQPtB5, float, 1.f / GPUCA_MIN_TRACK_PTB5_DEFAULT, "", 0, "required max Q/Pt (==min Pt) of tracks")
100100
AddOptionRTC(nonConsecutiveIDs, char, false, "", 0, "Non-consecutive cluster IDs as in HLT, disables features that need access to slice data in TPC merger")
101101
AddOptionRTC(fwdTPCDigitsAsClusters, unsigned char, 0, "", 0, "Forward TPC digits as clusters (if they pass the ZS threshold)")
102102
AddOptionRTC(bz0Pt10MeV, unsigned char, 60, "", 0, "Nominal Pt to set when bz = 0 (in 10 MeV)")
103103
AddOptionRTC(fitInProjections, char, -1, "", 0, "Fit in projection, -1 to enable for all but passes but the first one")
104104
AddOptionRTC(fitPropagateBzOnly, char, -1, "", 0, "Propagate using Bz only for n passes")
105105
AddOptionRTC(useMatLUT, char, 0, "", 0, "Use material lookup table for TPC refit")
106106
AddOptionRTC(trackingRefitGPUModel, char, 1, "", 0, "Use GPU track model for the Global Track Refit")
107-
AddCustomCPP(void SetMinTrackPt(float v) { maxTrackQPt = v > 0.001 ? (1. / v) : (1. / 0.001); })
107+
AddCustomCPP(void SetMinTrackPtB5(float v) { maxTrackQPtB5 = v > 0.001 ? (1. / v) : (1. / 0.001); })
108108
AddSubConfig(GPUSettingsRecTPC, tpc)
109109
AddSubConfig(GPUSettingsRecTRD, trd)
110110
AddHelp("help", 'h')
@@ -173,7 +173,6 @@ AddOption(clearO2OutputFromGPU, bool, false, "", 0, "Free the GPU memory used fo
173173
AddOption(ignoreNonFatalGPUErrors, bool, false, "", 0, "Continue running after having received non fatal GPU errors, e.g. abort due to overflow")
174174
AddOption(tpcIncreasedMinClustersPerRow, unsigned int, 0, "", 0, "Impose a minimum buffer size for the clustersPerRow during TPC clusterization")
175175
AddOption(noGPUMemoryRegistration, bool, false, "", 0, "Do not register input / output memory for GPU dma transfer")
176-
AddOption(automaticQPtThresholds, bool, true, "", 0, "Update the QPt thresholde at initialization accoding to the B field (i.e. lower to 40% for B=0.2T)")
177176
AddOption(calibObjectsExtraMemorySize, unsigned long, 10ul * 1024 * 1024, "", 0, "Extra spare memory added for calibration object buffer, to allow fow updates with larger objects")
178177
AddOption(useInternalO2Propagator, bool, false, "", 0, "Uses an internal (in GPUChainTracking) version of o2::Propagator, which internal b-field, matlut, etc.")
179178
AddOption(internalO2PropagatorGPUField, bool, true, "", 0, "Makes the internal O2 propagator use the fast GPU polynomial b field approximation")

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2008,12 +2008,12 @@ GPUd() void GPUTPCGMMerger::Finalize2(int nBlocks, int nThreads, int iBlock, int
20082008

20092009
GPUd() void GPUTPCGMMerger::MergeLoopersInit(int nBlocks, int nThreads, int iBlock, int iThread)
20102010
{
2011-
const float lowPtThresh = Param().rec.tpc.rejectQPt * 1.1f; // Might need to merge tracks above the threshold with parts below the threshold
2011+
const float lowPtThresh = Param().rec.tpc.rejectQPtB5 * 1.1f; // Might need to merge tracks above the threshold with parts below the threshold
20122012
for (unsigned int i = get_global_id(0); i < mMemory->nOutputTracks; i += get_global_size(0)) {
20132013
const auto& trk = mOutputTracks[i];
20142014
const auto& p = trk.GetParam();
20152015
const float qptabs = CAMath::Abs(p.GetQPt());
2016-
if (trk.NClusters() && qptabs * Param().par.qptB5Scaler > 5.f && qptabs <= lowPtThresh) {
2016+
if (trk.NClusters() && qptabs * Param().par.qptB5Scaler > 5.f && qptabs * Param().par.qptB5Scaler <= lowPtThresh) {
20172017
const int slice = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].slice;
20182018
const float refz = p.GetZ() + (Param().par.earlyTpcTransform ? p.GetTZOffset() : GetConstantMem()->calibObjects.fastTransform->convVertexTimeToZOffset(slice, p.GetTZOffset(), Param().par.continuousMaxTimeBin)) + (trk.CSide() ? -100 : 100);
20192019
float sinA, cosA;
@@ -2175,7 +2175,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int nBlocks, int nThreads, int iBlo
21752175
#endif
21762176
if (EQ) {
21772177
mOutputTracks[params[j].id].SetMergedLooper(true);
2178-
if (CAMath::Abs(param2.GetQPt()) >= Param().rec.tpc.rejectQPt) {
2178+
if (CAMath::Abs(param2.GetQPt() * Param().par.qptB5Scaler) >= Param().rec.tpc.rejectQPtB5) {
21792179
mOutputTracks[params[i].id].SetMergedLooper(true);
21802180
}
21812181
}

‎GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -58,7 +58,7 @@ MEM_CLASS_PRE23()
5858
GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int /*nBlocks*/, int /*nThreads*/, int /*iBlock*/, int /*iThread*/, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() MEM_LG2(GPUTPCTracker) & GPUrestrict() tracker, MEM_LG3(GPUTPCTrackParam) & GPUrestrict() tParam, calink* rowHits)
5959
{
6060
// reconstruction of tracklets, tracklet store step
61-
if (r.mNHits == 0 || (r.mNHits < GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(tParam.QPt() * tracker.Param().par.qptB5Scaler) || !CheckCov(tParam) || CAMath::Abs(tParam.GetQPt()) > tracker.Param().rec.maxTrackQPt)) {
61+
if (r.mNHits == 0 || (r.mNHits < GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(tParam.QPt() * tracker.Param().par.qptB5Scaler) || !CheckCov(tParam) || CAMath::Abs(tParam.GetQPt() * tracker.Param().par.qptB5Scaler) > tracker.Param().rec.maxTrackQPtB5)) {
6262
return;
6363
}
6464

‎GPU/GPUTracking/qa/GPUQA.cxx‎

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -88,20 +88,20 @@ using namespace GPUCA_NAMESPACE::gpu;
8888
#define CHECK_CLUSTER_STATE_INIT_LEG_BY_MC()
8989
#endif
9090

91-
#define CHECK_CLUSTER_STATE_INIT() \
92-
bool unattached = attach == 0; \
93-
float qpt = 0; \
94-
bool lowPt = false; \
95-
bool mev200 = false; \
96-
bool mergedLooper = false; \
97-
int id = attach & gputpcgmmergertypes::attachTrackMask; \
98-
if (!unattached) { \
99-
qpt = fabsf(mTracking->mIOPtrs.mergedTracks[id].GetParam().GetQPt()); \
100-
lowPt = qpt > mTracking->GetParam().rec.tpc.rejectQPt; \
101-
mev200 = qpt > 5; \
102-
mergedLooper = mTracking->mIOPtrs.mergedTracks[id].MergedLooper(); \
103-
} \
104-
bool physics = false, protect = false; \
91+
#define CHECK_CLUSTER_STATE_INIT() \
92+
bool unattached = attach == 0; \
93+
float qpt = 0; \
94+
bool lowPt = false; \
95+
bool mev200 = false; \
96+
bool mergedLooper = false; \
97+
int id = attach & gputpcgmmergertypes::attachTrackMask; \
98+
if (!unattached) { \
99+
qpt = fabsf(mTracking->mIOPtrs.mergedTracks[id].GetParam().GetQPt()); \
100+
lowPt = qpt * mTracking->GetParam().par.qptB5Scaler > mTracking->GetParam().rec.tpc.rejectQPtB5; \
101+
mev200 = qpt > 5; \
102+
mergedLooper = mTracking->mIOPtrs.mergedTracks[id].MergedLooper(); \
103+
} \
104+
bool physics = false, protect = false; \
105105
CHECK_CLUSTER_STATE_INIT_LEG_BY_MC();
106106

107107
#define CHECK_CLUSTER_STATE() \

‎prodtests/sim_challenge.sh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -127,7 +127,7 @@ if [ "$doreco" == "1" ]; then
127127

128128
echo "Running TPC reco flow"
129129
#needs TPC digitized data
130-
taskwrapper tpcreco.log o2-tpc-reco-workflow $gloOpt --input-type digits --output-type clusters,tracks,send-clusters-per-sector --configKeyValues "GPU_rec.maxTrackQPt=20"
130+
taskwrapper tpcreco.log o2-tpc-reco-workflow $gloOpt --input-type digits --output-type clusters,tracks,send-clusters-per-sector --configKeyValues "GPU_rec.maxTrackQPtB5=20"
131131
echo "Return status of tpcreco: $?"
132132

133133
echo "Running ITS reco flow"

0 commit comments

Comments
 (0)