Skip to content

Commit 0815f58

Browse files
f3schf3sch
committed
more refactor
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 75c79fe commit 0815f58

3 files changed

Lines changed: 145 additions & 165 deletions

File tree

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu‎

Lines changed: 22 additions & 81 deletions
Original file line numberDiff line numberDiff line change
@@ -134,7 +134,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel(
134134
const float maxChi2NDF,
135135
const int reseedIfShorter,
136136
const bool repeatRefitOut,
137-
const bool shifRefToCluster,
137+
const bool shiftRefToCluster,
138138
const o2::base::Propagator* propagator,
139139
const o2::base::PropagatorF::MatCorrType matCorrType)
140140
{
@@ -145,87 +145,28 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel(
145145
continue;
146146
}
147147
}
148-
149-
TrackITSExt temporaryTrack = o2::its::track::seedTrackForRefit(trackSeeds[iCurrentTrackSeedIndex],
150-
foundTrackingFrameInfo,
151-
unsortedClusters,
152-
layerRadii,
153-
bz,
154-
reseedIfShorter);
155-
o2::track::TrackPar linRef{temporaryTrack};
156-
bool fitSuccess = o2::its::track::fitTrack(temporaryTrack, // TrackITSExt& track,
157-
0, // int lastLayer,
158-
NLayers, // int firstLayer,
159-
1, // int firstCluster,
160-
maxChi2ClusterAttachment, // float maxChi2ClusterAttachment,
161-
maxChi2NDF, // float maxChi2NDF,
162-
o2::constants::math::VeryBig, // float maxQoverPt,
163-
0, // nCl,
164-
bz, // float bz,
165-
foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo,
166-
layerxX0, // const float* layerxX0,
167-
propagator, // const o2::base::Propagator* propagator,
168-
matCorrType, // o2::base::PropagatorF::MatCorrType matCorrType
169-
&linRef,
170-
shifRefToCluster);
171-
if (!fitSuccess) {
172-
continue;
173-
}
174-
temporaryTrack.getParamOut() = temporaryTrack.getParamIn();
175-
linRef = temporaryTrack.getParamOut(); // use refitted track as lin.reference
176-
track::resetTrackCovariance(temporaryTrack);
177-
temporaryTrack.setChi2(0);
178-
fitSuccess = o2::its::track::fitTrack(temporaryTrack, // TrackITSExt& track,
179-
NLayers - 1, // int lastLayer,
180-
-1, // int firstLayer,
181-
-1, // int firstCluster,
182-
maxChi2ClusterAttachment, // float maxChi2ClusterAttachment,
183-
maxChi2NDF, // float maxChi2NDF,
184-
50.f, // float maxQoverPt,
185-
0, // nCl,
186-
bz, // float bz,
187-
foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo,
188-
layerxX0, // const float* layerxX0,
189-
propagator, // const o2::base::Propagator* propagator,
190-
matCorrType, // o2::base::PropagatorF::MatCorrType matCorrType
191-
&linRef,
192-
shifRefToCluster);
193-
if (!fitSuccess || temporaryTrack.getPt() < minPts[NLayers - temporaryTrack.getNClusters()]) {
194-
continue;
195-
}
196-
if (repeatRefitOut) { // repeat outward refit seeding and linearizing with the stable inward fit result
197-
o2::track::TrackParCov saveInw{temporaryTrack};
198-
linRef = saveInw; // use refitted track as lin.reference
199-
float saveChi2 = temporaryTrack.getChi2();
200-
track::resetTrackCovariance(temporaryTrack);
201-
temporaryTrack.setChi2(0);
202-
fitSuccess = o2::its::track::fitTrack(temporaryTrack, // TrackITSExt& track,
203-
0, // int lastLayer,
204-
NLayers, // int firstLayer,
205-
1, // int firstCluster,
206-
maxChi2ClusterAttachment, // float maxChi2ClusterAttachment,
207-
maxChi2NDF, // float maxChi2NDF,
208-
o2::constants::math::VeryBig, // float maxQoverPt,
209-
0, // nCl,
210-
bz, // float bz,
211-
foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo,
212-
layerxX0, // const float* layerxX0,
213-
propagator, // const o2::base::Propagator* propagator,
214-
matCorrType, // o2::base::PropagatorF::MatCorrType matCorrType
215-
&linRef,
216-
shifRefToCluster);
217-
if (!fitSuccess) {
218-
continue;
148+
TrackITSExt temporaryTrack;
149+
bool refitSuccess = o2::its::track::refitTrack(trackSeeds[iCurrentTrackSeedIndex],
150+
temporaryTrack,
151+
maxChi2ClusterAttachment,
152+
maxChi2NDF,
153+
bz,
154+
foundTrackingFrameInfo,
155+
unsortedClusters,
156+
layerxX0,
157+
layerRadii,
158+
minPts,
159+
propagator,
160+
matCorrType,
161+
reseedIfShorter,
162+
shiftRefToCluster,
163+
repeatRefitOut);
164+
if (refitSuccess) {
165+
if constexpr (initRun) {
166+
seedLUT[iCurrentTrackSeedIndex] = 1;
167+
} else {
168+
tracks[seedLUT[iCurrentTrackSeedIndex]] = temporaryTrack;
219169
}
220-
temporaryTrack.getParamOut() = temporaryTrack.getParamIn();
221-
temporaryTrack.getParamIn() = saveInw;
222-
temporaryTrack.setChi2(saveChi2);
223-
}
224-
225-
if constexpr (initRun) {
226-
seedLUT[iCurrentTrackSeedIndex] = 1;
227-
} else {
228-
tracks[seedLUT[iCurrentTrackSeedIndex]] = temporaryTrack;
229170
}
230171
}
231172
}

‎Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h‎

Lines changed: 95 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -191,6 +191,101 @@ GPUdi() bool fitTrack(TrackITSExt& trk,
191191
return o2::gpu::CAMath::Abs(trk.getQ2Pt()) < maxQoverPt && trk.getChi2() < chi2ndfcut * (float)((nCl * 2) - 5);
192192
}
193193

194+
template <int NLayers>
195+
GPUdi() bool refitTrack(const TrackSeed<NLayers>& trackSeed,
196+
TrackITSExt& temporaryTrack,
197+
float chi2clcut,
198+
float chi2ndfcut,
199+
const float bz,
200+
const TrackingFrameInfo* const* tfInfos,
201+
const Cluster* const* clusters,
202+
const float* layerxX0,
203+
const float* layerRadii,
204+
const float* minPt,
205+
const o2::base::Propagator* propagator,
206+
const o2::base::PropagatorF::MatCorrType matCorrType,
207+
const int reseedIfShorter,
208+
const bool shiftRefToCluster,
209+
const bool repeatRefitOut)
210+
{
211+
temporaryTrack = seedTrackForRefit(trackSeed,
212+
tfInfos,
213+
clusters,
214+
layerRadii,
215+
bz,
216+
reseedIfShorter);
217+
o2::track::TrackPar linRef{temporaryTrack};
218+
bool fitSuccess = fitTrack(temporaryTrack,
219+
0,
220+
NLayers,
221+
1,
222+
chi2clcut,
223+
chi2ndfcut,
224+
o2::constants::math::VeryBig,
225+
0,
226+
bz,
227+
tfInfos,
228+
layerxX0,
229+
propagator,
230+
matCorrType,
231+
&linRef,
232+
shiftRefToCluster);
233+
if (!fitSuccess) {
234+
return false;
235+
}
236+
temporaryTrack.getParamOut() = temporaryTrack.getParamIn();
237+
linRef = temporaryTrack.getParamOut(); // use refitted track as lin.reference
238+
resetTrackCovariance(temporaryTrack);
239+
temporaryTrack.setChi2(0);
240+
fitSuccess = fitTrack(temporaryTrack,
241+
NLayers - 1,
242+
-1,
243+
-1,
244+
chi2clcut,
245+
chi2ndfcut,
246+
50.f,
247+
0,
248+
bz,
249+
tfInfos,
250+
layerxX0,
251+
propagator,
252+
matCorrType,
253+
&linRef,
254+
shiftRefToCluster);
255+
if (!fitSuccess || temporaryTrack.getPt() < minPt[NLayers - temporaryTrack.getNClusters()]) {
256+
return false;
257+
}
258+
if (repeatRefitOut) { // repeat outward refit seeding and linearizing with the stable inward fit result
259+
o2::track::TrackParCov saveInw{temporaryTrack};
260+
linRef = saveInw; // use refitted track as lin.reference
261+
float saveChi2 = temporaryTrack.getChi2();
262+
track::resetTrackCovariance(temporaryTrack);
263+
temporaryTrack.setChi2(0);
264+
fitSuccess = o2::its::track::fitTrack(temporaryTrack,
265+
0,
266+
NLayers,
267+
1,
268+
chi2clcut,
269+
chi2ndfcut,
270+
o2::constants::math::VeryBig,
271+
0,
272+
bz,
273+
tfInfos,
274+
layerxX0,
275+
propagator,
276+
matCorrType,
277+
&linRef,
278+
shiftRefToCluster);
279+
if (!fitSuccess) {
280+
return false;
281+
}
282+
temporaryTrack.getParamOut() = temporaryTrack.getParamIn();
283+
temporaryTrack.getParamIn() = saveInw;
284+
temporaryTrack.setChi2(saveChi2);
285+
}
286+
return true;
287+
}
288+
194289
} // namespace o2::its::track
195290

196291
#endif

‎Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx‎

Lines changed: 28 additions & 84 deletions
Original file line numberDiff line numberDiff line change
@@ -713,92 +713,36 @@ void TrackerTraits<NLayers>::findRoads(const int iteration)
713713
bounded_vector<TrackITSExt> tracks(mMemoryPool.get());
714714
mTaskArena->execute([&] {
715715
auto forSeed = [&](auto Tag, int iSeed, int offset = 0) {
716-
TrackITSExt temporaryTrack = o2::its::track::seedTrackForRefit(trackSeeds[iSeed],
717-
tfInfos,
718-
unsortedClusters,
719-
mTrkParams[iteration].LayerRadii.data(),
720-
mBz,
721-
mTrkParams[iteration].ReseedIfShorter);
722-
o2::track::TrackPar linRef{temporaryTrack};
723-
bool fitSuccess = o2::its::track::fitTrack(temporaryTrack,
724-
0,
725-
mTrkParams[iteration].NLayers,
726-
1,
727-
mTrkParams[iteration].MaxChi2ClusterAttachment,
728-
mTrkParams[iteration].MaxChi2NDF,
729-
o2::constants::math::VeryBig,
730-
0,
731-
mBz,
732-
tfInfos,
733-
mTrkParams[iteration].LayerxX0.data(),
734-
propagator,
735-
mTrkParams[iteration].CorrType,
736-
&linRef,
737-
mTrkParams[iteration].ShiftRefToCluster);
738-
if (!fitSuccess) {
739-
return 0;
740-
}
741-
temporaryTrack.getParamOut() = temporaryTrack.getParamIn();
742-
linRef = temporaryTrack.getParamOut(); // use refitted track as lin.reference
743-
track::resetTrackCovariance(temporaryTrack);
744-
temporaryTrack.setChi2(0);
745-
fitSuccess = o2::its::track::fitTrack(temporaryTrack,
746-
mTrkParams[iteration].NLayers - 1,
747-
-1,
748-
-1,
749-
mTrkParams[iteration].MaxChi2ClusterAttachment,
750-
mTrkParams[iteration].MaxChi2NDF,
751-
50.f,
752-
0,
753-
mBz,
754-
tfInfos,
755-
mTrkParams[iteration].LayerxX0.data(),
756-
propagator,
757-
mTrkParams[iteration].CorrType,
758-
&linRef,
759-
mTrkParams[iteration].ShiftRefToCluster);
760-
if (!fitSuccess || temporaryTrack.getPt() < mTrkParams[iteration].MinPt[mTrkParams[iteration].NLayers - temporaryTrack.getNClusters()]) {
761-
return 0;
762-
}
763-
if (mTrkParams[iteration].RepeatRefitOut) { // repeat outward refit seeding and linearizing with the stable inward fit result
764-
o2::track::TrackParCov saveInw{temporaryTrack};
765-
linRef = saveInw; // use refitted track as lin.reference
766-
float saveChi2 = temporaryTrack.getChi2();
767-
track::resetTrackCovariance(temporaryTrack);
768-
temporaryTrack.setChi2(0);
769-
fitSuccess = o2::its::track::fitTrack(temporaryTrack,
770-
0,
771-
mTrkParams[iteration].NLayers,
772-
1,
773-
mTrkParams[iteration].MaxChi2ClusterAttachment,
774-
mTrkParams[iteration].MaxChi2NDF,
775-
o2::constants::math::VeryBig,
776-
0,
777-
mBz,
778-
tfInfos,
779-
mTrkParams[iteration].LayerxX0.data(),
780-
propagator,
781-
mTrkParams[iteration].CorrType,
782-
&linRef,
783-
mTrkParams[iteration].ShiftRefToCluster);
784-
if (!fitSuccess) {
785-
return 0;
716+
TrackITSExt temporaryTrack;
717+
bool refitSuccess = track::refitTrack<NLayers>(trackSeeds[iSeed],
718+
temporaryTrack,
719+
mTrkParams[iteration].MaxChi2ClusterAttachment,
720+
mTrkParams[iteration].MaxChi2NDF,
721+
mBz,
722+
tfInfos,
723+
unsortedClusters,
724+
mTrkParams[iteration].LayerxX0.data(),
725+
mTrkParams[iteration].LayerRadii.data(),
726+
mTrkParams[iteration].MinPt.data(),
727+
propagator,
728+
mTrkParams[iteration].CorrType,
729+
mTrkParams[iteration].ReseedIfShorter,
730+
mTrkParams[iteration].ShiftRefToCluster,
731+
mTrkParams[iteration].RepeatRefitOut);
732+
733+
if (refitSuccess) {
734+
if constexpr (decltype(Tag)::value == PassMode::OnePass::value) {
735+
tracks.push_back(temporaryTrack);
736+
} else if constexpr (decltype(Tag)::value == PassMode::TwoPassCount::value) {
737+
// nothing to do
738+
} else if constexpr (decltype(Tag)::value == PassMode::TwoPassInsert::value) {
739+
tracks[offset] = temporaryTrack;
740+
} else {
741+
static_assert(false, "Unknown mode!");
786742
}
787-
temporaryTrack.getParamOut() = temporaryTrack.getParamIn();
788-
temporaryTrack.getParamIn() = saveInw;
789-
temporaryTrack.setChi2(saveChi2);
743+
return 1;
790744
}
791-
792-
if constexpr (decltype(Tag)::value == PassMode::OnePass::value) {
793-
tracks.push_back(temporaryTrack);
794-
} else if constexpr (decltype(Tag)::value == PassMode::TwoPassCount::value) {
795-
// nothing to do
796-
} else if constexpr (decltype(Tag)::value == PassMode::TwoPassInsert::value) {
797-
tracks[offset] = temporaryTrack;
798-
} else {
799-
static_assert(false, "Unknown mode!");
800-
}
801-
return 1;
745+
return 0;
802746
};
803747

804748
const int nSeeds = static_cast<int>(trackSeeds.size());

0 commit comments

Comments
 (0)