Skip to content

Commit d0bd615

Browse files
mpucciompuccio
committed
ITS: simplifications
1 parent 507c544 commit d0bd615

9 files changed

Lines changed: 67 additions & 93 deletions

File tree

‎Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -97,7 +97,6 @@ class TimeFrameGPU : public TimeFrame<NLayers>
9797
void createTrackExtensionResultsDevice(const size_t);
9898
void downloadTrackITSExtDevice();
9999
void downloadCellsNeighboursDevice(std::vector<bounded_vector<CellNeighbour>>&, const int);
100-
void downloadTrackExtensionCandidatesDevice();
101100
void downloadTrackExtensionResultsDevice();
102101
void downloadNeighboursLUTDevice(bounded_vector<int>&, const int);
103102
void downloadCellsDevice();
@@ -125,7 +124,6 @@ class TimeFrameGPU : public TimeFrame<NLayers>
125124
const auto getDeviceTrackingTopologyView() const { return mDeviceTrackingTopologyView; }
126125
int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; }
127126
auto& getTrackITSExt() { return mTrackITSExt; }
128-
auto& getTrackExtensionCandidates() { return mTrackExtensionCandidates; }
129127
auto& getTrackExtensionResults() { return mTrackExtensionResults; }
130128
Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; }
131129
int* getDeviceROFramesPV() { return mROFramesPVDevice; }
@@ -261,8 +259,6 @@ class TimeFrameGPU : public TimeFrame<NLayers>
261259
// Temporary buffer for storing output tracks from GPU tracking
262260
bounded_vector<TrackITSExt> mTrackITSExt;
263261
bounded_vector<TrackITSExt> mTrackExtensionStartTracks;
264-
// Temporary buffer for compact track extension proposals from GPU tracking
265-
bounded_vector<TrackExtensionCandidate<NLayers>> mTrackExtensionCandidates;
266262
// Temporary buffer for fitted track extension proposals from GPU tracking
267263
bounded_vector<TrackExtensionResult<NLayers>> mTrackExtensionResults;
268264
};

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

Lines changed: 1 addition & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -602,10 +602,9 @@ void TimeFrameGPU<NLayers>::createTrackExtensionCandidatesDevice(const size_t nT
602602
GPUTimer timer("reserving track extension candidates");
603603
const size_t nCandidates = nTracks * MaxTrackExtensionCandidatesPerTrack;
604604
GPULog("gpu-allocation: reserving {} track extension candidates, for {:.2f} MB.", nCandidates, nCandidates * sizeof(o2::its::TrackExtensionCandidate<NLayers>) / constants::MB);
605-
mTrackExtensionCandidates = bounded_vector<TrackExtensionCandidate<NLayers>>(nCandidates, {}, this->getMemoryPool().get());
606605
mTrackExtensionCandidatesDevice = nullptr;
607606
mTrackExtensionCandidateOffsetsDevice = nullptr;
608-
if (mTrackExtensionCandidates.empty()) {
607+
if (nCandidates == 0) {
609608
return;
610609
}
611610
allocMem(reinterpret_cast<void**>(&mTrackExtensionCandidatesDevice), nCandidates * sizeof(o2::its::TrackExtensionCandidate<NLayers>), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
@@ -695,17 +694,6 @@ void TimeFrameGPU<NLayers>::downloadTrackITSExtDevice()
695694
GPUChkErrS(cudaMemcpy(mTrackITSExt.data(), mTrackITSExtDevice, mTrackITSExt.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyDeviceToHost));
696695
}
697696

698-
template <int NLayers>
699-
void TimeFrameGPU<NLayers>::downloadTrackExtensionCandidatesDevice()
700-
{
701-
GPUTimer timer("downloading track extension candidates");
702-
GPULog("gpu-transfer: downloading {} track extension candidates, for {:.2f} MB.", mTrackExtensionCandidates.size(), mTrackExtensionCandidates.size() * sizeof(o2::its::TrackExtensionCandidate<NLayers>) / constants::MB);
703-
if (mTrackExtensionCandidates.empty()) {
704-
return;
705-
}
706-
GPUChkErrS(cudaMemcpy(mTrackExtensionCandidates.data(), mTrackExtensionCandidatesDevice, mTrackExtensionCandidates.size() * sizeof(o2::its::TrackExtensionCandidate<NLayers>), cudaMemcpyDeviceToHost));
707-
}
708-
709697
template <int NLayers>
710698
void TimeFrameGPU<NLayers>::downloadTrackExtensionResultsDevice()
711699
{

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -485,20 +485,20 @@ void TrackerTraitsGPU<NLayers>::buildTrackExtensionCandidates(const int iteratio
485485
continue;
486486
}
487487
auto candidate = result.candidate;
488-
candidate.fittedTrackIndex = iResult;
489-
candidatesPerTrack[result.trackIndex].push_back(candidate);
488+
candidate.resultIndex = iResult;
489+
candidatesPerTrack[candidate.trackIndex].push_back(candidate);
490490
}
491491
}
492492

493493
template <int NLayers>
494494
bool TrackerTraitsGPU<NLayers>::materializeTrackExtensionCandidate(TrackITSExt& track, const typename TrackerTraits<NLayers>::TrackExtensionCandidateN& candidate, const int iteration)
495495
{
496496
const auto& results = mTimeFrameGPU->getTrackExtensionResults();
497-
if (candidate.fittedTrackIndex < 0 || candidate.fittedTrackIndex >= static_cast<int>(results.size())) {
497+
if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast<int>(results.size())) {
498498
return TrackerTraits<NLayers>::materializeTrackExtensionCandidate(track, candidate, iteration);
499499
}
500-
const auto& result = results[candidate.fittedTrackIndex];
501-
if (!result.isValid() || result.trackIndex != candidate.trackIndex) {
500+
const auto& result = results[candidate.resultIndex];
501+
if (!result.isValid() || result.candidate.trackIndex != candidate.trackIndex) {
502502
return false;
503503
}
504504
track = result.track;

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

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -134,7 +134,6 @@ GPUdi() void writeTrackExtensionCandidate(const int trackIndex,
134134
return;
135135
}
136136
candidate.chi2 = updated.getChi2();
137-
candidate.chi2NDF = updated.getChi2() / static_cast<float>(updated.getNClusters() * 2 - 5);
138137
++slot;
139138
}
140139

@@ -313,7 +312,6 @@ GPUg() void __launch_bounds__(256, 1) computeTrackExtensionResultsKernel(const T
313312
if (!candidate.isValidForTrack(iTrack)) {
314313
continue;
315314
}
316-
result.trackIndex = iTrack;
317315
result.candidate = candidate;
318316
if (!fitTrackExtensionResult(startTrack,
319317
candidate,
@@ -331,7 +329,6 @@ GPUg() void __launch_bounds__(256, 1) computeTrackExtensionResultsKernel(const T
331329
continue;
332330
}
333331
result.candidate.chi2 = result.track.getChi2();
334-
result.candidate.chi2NDF = result.track.getChi2() / static_cast<float>(result.track.getNClusters() * 2 - 5);
335332
}
336333
}
337334
}

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

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -93,6 +93,45 @@ struct LayerTiming {
9393
return (bc - offset) / mROFLength;
9494
}
9595

96+
// return which ROF this floating point (number of BCs) time belongs
97+
GPUhdi() BCType getROF(float time) const noexcept
98+
{
99+
const float offset = static_cast<float>(mROFDelay + mROFBias);
100+
if (time <= offset) {
101+
return 0;
102+
}
103+
return static_cast<BCType>((time - offset) / mROFLength);
104+
}
105+
106+
GPUhdi() bool intersectROF(BCType rof, float lower, float upper) const noexcept
107+
{
108+
const auto rofTS = getROFTimeBounds(rof, true);
109+
return static_cast<float>(rofTS.upper()) > lower && upper > static_cast<float>(rofTS.lower());
110+
}
111+
112+
// return clamped ROF range with strictly positive overlap with timestamp interval
113+
GPUhdi() int2 getROFRange(TimeStamp ts) const noexcept
114+
{
115+
if (mNROFsTF == 0) {
116+
return {1, 0};
117+
}
118+
119+
const float lower = ts.getTimeStamp() - ts.getTimeStampError();
120+
const float upper = ts.getTimeStamp() + ts.getTimeStampError();
121+
const int maxROF = static_cast<int>(mNROFsTF) - 1;
122+
int2 range{
123+
o2::gpu::CAMath::Clamp(static_cast<int>(getROF(lower - mROFAddTimeErr)), 0, maxROF),
124+
o2::gpu::CAMath::Clamp(static_cast<int>(getROF(upper + mROFAddTimeErr)), 0, maxROF)};
125+
126+
if (range.x <= range.y && !intersectROF(static_cast<BCType>(range.x), lower, upper)) {
127+
++range.x;
128+
}
129+
if (range.y >= range.x && !intersectROF(static_cast<BCType>(range.y), lower, upper)) {
130+
--range.y;
131+
}
132+
return range;
133+
}
134+
96135
#ifndef GPUCA_GPUCODE
97136
GPUh() std::string asString() const
98137
{

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -279,7 +279,7 @@ struct TimeFrame {
279279
std::vector<bounded_vector<Tracklet>> mTracklets;
280280
std::vector<bounded_vector<CellSeed>> mCells;
281281
bounded_vector<TrackITSExt> mTracks;
282-
std::vector<std::vector<TrackITSExt>> mFittedExtensionTracks;
282+
std::vector<TrackITSExt> mFittedExtensionTracks;
283283
bounded_vector<MCCompLabel> mTracksLabel;
284284
std::vector<bounded_vector<int>> mCellsNeighbours;
285285
std::vector<bounded_vector<int>> mCellsNeighboursTopology;

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

Lines changed: 9 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -64,8 +64,7 @@ struct TrackExtensionCandidate {
6464
{
6565
trackIndex = -1;
6666
nAddedClusters = 0;
67-
fittedTrackIndex = -1;
68-
chi2NDF = InvalidChi2;
67+
resultIndex = -1;
6968
chi2 = InvalidChi2;
7069
for (int iLayer{0}; iLayer < NLayers; ++iLayer) {
7170
addedClusters[iLayer] = constants::UnusedIndex;
@@ -80,22 +79,25 @@ struct TrackExtensionCandidate {
8079
int trackIndex{-1};
8180
std::array<int, NLayers> addedClusters;
8281
int nAddedClusters{0};
83-
int fittedTrackIndex{-1};
84-
float chi2NDF{InvalidChi2};
82+
int resultIndex{-1};
8583
float chi2{InvalidChi2};
8684
};
8785

86+
template <int NLayers>
87+
GPUhdi() bool isBetterTrackExtensionCandidate(const TrackExtensionCandidate<NLayers>& a, const TrackExtensionCandidate<NLayers>& b)
88+
{
89+
return (a.nAddedClusters > b.nAddedClusters) || (a.nAddedClusters == b.nAddedClusters && a.chi2 < b.chi2);
90+
}
91+
8892
template <int NLayers>
8993
struct TrackExtensionResult {
9094
GPUhdi() void reset()
9195
{
92-
trackIndex = -1;
9396
candidate.reset();
9497
}
9598

96-
GPUhdi() bool isValid() const { return trackIndex >= 0 && candidate.nAddedClusters > 0; }
99+
GPUhdi() bool isValid() const { return candidate.trackIndex >= 0 && candidate.nAddedClusters > 0; }
97100

98-
int trackIndex{-1};
99101
TrackExtensionCandidate<NLayers> candidate;
100102
TrackITSExt track;
101103
};

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

Lines changed: 3 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -33,13 +33,7 @@ namespace o2::its
3333
template <int NLayers>
3434
GPUhdi() bool isBetterTrackExtensionHypothesis(const TrackExtensionHypothesis<NLayers>& a, const TrackExtensionHypothesis<NLayers>& b)
3535
{
36-
if (a.nClusters != b.nClusters) {
37-
return a.nClusters > b.nClusters;
38-
}
39-
if (a.chi2 != b.chi2) {
40-
return a.chi2 < b.chi2;
41-
}
42-
return false;
36+
return (a.nClusters > b.nClusters) || (a.nClusters == b.nClusters && a.chi2 < b.chi2);
4337
}
4438

4539
template <int NLayers>
@@ -85,31 +79,6 @@ GPUhdi() int4 getTrackExtensionBinsAt(const IndexTableUtils<NLayers>& utils,
8579
utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))};
8680
}
8781

88-
template <int NLayers>
89-
GPUhdi() int2 getTrackExtensionCompatibleROFRange(const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
90-
const int layer,
91-
const TimeStamp& time)
92-
{
93-
const auto& timing = rofOverlaps.getLayer(layer);
94-
int first = static_cast<int>(timing.getROF(static_cast<LayerTiming::BCType>(o2::gpu::CAMath::Max(time.getTimeStamp() - time.getTimeStampError(), 0.f))));
95-
int last = static_cast<int>(timing.getROF(static_cast<LayerTiming::BCType>(o2::gpu::CAMath::Max(time.getTimeStamp() + time.getTimeStampError(), 0.f))));
96-
first = o2::gpu::CAMath::Min(o2::gpu::CAMath::Max(first, 0), static_cast<int>(timing.mNROFsTF) - 1);
97-
last = o2::gpu::CAMath::Min(o2::gpu::CAMath::Max(last, 0), static_cast<int>(timing.mNROFsTF) - 1);
98-
return {first, last};
99-
}
100-
101-
template <int NLayers>
102-
GPUhdi() bool isTrackExtensionROFTimeCompatible(const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
103-
const int layer,
104-
const int rof,
105-
const TimeStamp& time)
106-
{
107-
const auto rofTS = rofOverlaps.getLayer(layer).getROFTimeBounds(rof, true);
108-
const float lower = time.getTimeStamp() - time.getTimeStampError();
109-
const float upper = time.getTimeStamp() + time.getTimeStampError();
110-
return static_cast<float>(rofTS.upper()) > lower && upper > static_cast<float>(rofTS.lower());
111-
}
112-
11382
template <int NLayers>
11483
GPUhdi() int getTrackExtensionFirstClusterLayer(const TrackITSExt& track)
11584
{
@@ -223,9 +192,9 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track,
223192
phiBinsNum += phiBins;
224193
}
225194

226-
const auto rofRange = getTrackExtensionCompatibleROFRange<NLayers>(rofOverlaps, iLayer, hypo.time);
195+
const auto rofRange = rofOverlaps.getLayer(iLayer).getROFRange(hypo.time);
227196
for (int rof = rofRange.x; rof <= rofRange.y; ++rof) {
228-
if (!rofMask.isROFEnabled(iLayer, rof) || !isTrackExtensionROFTimeCompatible<NLayers>(rofOverlaps, iLayer, rof, hypo.time)) {
197+
if (!rofMask.isROFEnabled(iLayer, rof)) {
229198
continue;
230199
}
231200
const int rofStart = ROFClusters[iLayer][rof];

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

Lines changed: 9 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -904,7 +904,8 @@ void TrackerTraits<NLayers>::extendTracks(const int iteration)
904904
{
905905
const auto nTracks = mTimeFrame->getTracks().size();
906906
TrackExtensionCandidates candidatesPerTrack(nTracks);
907-
mTimeFrame->mFittedExtensionTracks.assign(nTracks, {});
907+
mTimeFrame->mFittedExtensionTracks.clear();
908+
mTimeFrame->mFittedExtensionTracks.reserve(nTracks * MaxTrackExtensionCandidatesPerTrack);
908909
buildTrackExtensionCandidates(iteration, candidatesPerTrack);
909910
applyTrackExtensionCandidates(iteration, candidatesPerTrack);
910911
mTimeFrame->mFittedExtensionTracks.clear();
@@ -1000,14 +1001,10 @@ void TrackerTraits<NLayers>::updateExtendedTrackTimeStamp(TrackITSExt& track, co
10001001
template <int NLayers>
10011002
bool TrackerTraits<NLayers>::materializeTrackExtensionCandidate(TrackITSExt& track, const TrackExtensionCandidateN& candidate, const int /*iteration*/)
10021003
{
1003-
if (candidate.trackIndex < 0 || candidate.trackIndex >= static_cast<int>(mTimeFrame->mFittedExtensionTracks.size())) {
1004-
return false;
1005-
}
1006-
const auto& slot = mTimeFrame->mFittedExtensionTracks[candidate.trackIndex];
1007-
if (candidate.fittedTrackIndex < 0 || candidate.fittedTrackIndex >= static_cast<int>(slot.size())) {
1004+
if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast<int>(mTimeFrame->mFittedExtensionTracks.size())) {
10081005
return false;
10091006
}
1010-
track = slot[candidate.fittedTrackIndex];
1007+
track = mTimeFrame->mFittedExtensionTracks[candidate.resultIndex];
10111008
return true;
10121009
}
10131010

@@ -1037,11 +1034,9 @@ void TrackerTraits<NLayers>::buildTrackExtensionCandidates(const int iteration,
10371034
if (!extension.nAddedClusters) {
10381035
return;
10391036
}
1040-
extension.chi2NDF = candidate.getChi2() / static_cast<float>(candidate.getNClusters() * 2 - 5);
10411037
extension.chi2 = candidate.getChi2();
1042-
auto& fittedSlot = mTimeFrame->mFittedExtensionTracks[trackIndex];
1043-
extension.fittedTrackIndex = static_cast<int>(fittedSlot.size());
1044-
fittedSlot.push_back(candidate);
1038+
extension.resultIndex = static_cast<int>(mTimeFrame->mFittedExtensionTracks.size());
1039+
mTimeFrame->mFittedExtensionTracks.push_back(candidate);
10451040
candidates.push_back(extension);
10461041
};
10471042

@@ -1105,21 +1100,9 @@ template <int NLayers>
11051100
void TrackerTraits<NLayers>::applyTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack)
11061101
{
11071102
auto& tracks = mTimeFrame->getTracks();
1108-
auto isBetterCandidate = [](const TrackExtensionCandidateN& a, const TrackExtensionCandidateN& b) {
1109-
if (a.nAddedClusters != b.nAddedClusters) {
1110-
return a.nAddedClusters > b.nAddedClusters;
1111-
}
1112-
if (a.chi2NDF != b.chi2NDF) {
1113-
return a.chi2NDF < b.chi2NDF;
1114-
}
1115-
if (a.chi2 != b.chi2) {
1116-
return a.chi2 < b.chi2;
1117-
}
1118-
return false;
1119-
};
11201103

11211104
for (auto& candidates : candidatesPerTrack) {
1122-
std::stable_sort(candidates.begin(), candidates.end(), isBetterCandidate);
1105+
std::stable_sort(candidates.begin(), candidates.end(), isBetterTrackExtensionCandidate<NLayers>);
11231106
while (!candidates.empty() && (candidates.back().nAddedClusters <= 0)) {
11241107
candidates.pop_back();
11251108
}
@@ -1137,10 +1120,10 @@ void TrackerTraits<NLayers>::applyTrackExtensionCandidates(const int iteration,
11371120
auto cmp = [&](const Entry& a, const Entry& b) {
11381121
const auto& ca = candidatesPerTrack[a.track][a.idx];
11391122
const auto& cb = candidatesPerTrack[b.track][b.idx];
1140-
if (isBetterCandidate(cb, ca)) {
1123+
if (isBetterTrackExtensionCandidate<NLayers>(cb, ca)) {
11411124
return true;
11421125
}
1143-
if (isBetterCandidate(ca, cb)) {
1126+
if (isBetterTrackExtensionCandidate<NLayers>(ca, cb)) {
11441127
return false;
11451128
}
11461129
if (a.track != b.track) {

0 commit comments

Comments
 (0)