Skip to content

Commit 1cb33e2

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

11 files changed

Lines changed: 139 additions & 111 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: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -485,20 +485,22 @@ 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+
if (candidatesPerTrack.add(candidate.trackIndex, candidate) < 0) {
490+
continue;
491+
}
490492
}
491493
}
492494

493495
template <int NLayers>
494496
bool TrackerTraitsGPU<NLayers>::materializeTrackExtensionCandidate(TrackITSExt& track, const typename TrackerTraits<NLayers>::TrackExtensionCandidateN& candidate, const int iteration)
495497
{
496498
const auto& results = mTimeFrameGPU->getTrackExtensionResults();
497-
if (candidate.fittedTrackIndex < 0 || candidate.fittedTrackIndex >= static_cast<int>(results.size())) {
499+
if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast<int>(results.size())) {
498500
return TrackerTraits<NLayers>::materializeTrackExtensionCandidate(track, candidate, iteration);
499501
}
500-
const auto& result = results[candidate.fittedTrackIndex];
501-
if (!result.isValid() || result.trackIndex != candidate.trackIndex) {
502+
const auto& result = results[candidate.resultIndex];
503+
if (!result.isValid() || result.candidate.trackIndex != candidate.trackIndex) {
502504
return false;
503505
}
504506
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/include/ITStracking/TrackerTraits.h‎

Lines changed: 35 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,41 @@ class TrackerTraits
9494

9595
protected:
9696
using TrackExtensionCandidateN = TrackExtensionCandidate<NLayers>;
97-
using TrackExtensionCandidates = std::vector<std::vector<TrackExtensionCandidateN>>;
97+
struct TrackExtensionCandidates {
98+
TrackExtensionCandidates() = default;
99+
explicit TrackExtensionCandidates(size_t nTracks)
100+
: candidates(nTracks * MaxTrackExtensionCandidatesPerTrack), counts(nTracks, 0)
101+
{
102+
}
103+
104+
int add(int trackIndex, const TrackExtensionCandidateN& candidate)
105+
{
106+
auto& count = counts[trackIndex];
107+
if (count >= MaxTrackExtensionCandidatesPerTrack) {
108+
return -1;
109+
}
110+
const int flatIndex = static_cast<int>(getFlatTrackExtensionCandidateIndex(trackIndex, count));
111+
candidates[flatIndex] = candidate;
112+
++count;
113+
return flatIndex;
114+
}
115+
116+
void pop_back(int trackIndex)
117+
{
118+
--counts[trackIndex];
119+
}
120+
121+
bool empty(int trackIndex) const { return counts[trackIndex] == 0; }
122+
int size(int trackIndex) const { return counts[trackIndex]; }
123+
TrackExtensionCandidateN* begin(int trackIndex) { return candidates.data() + getFlatTrackExtensionCandidateIndex(trackIndex, 0); }
124+
TrackExtensionCandidateN* end(int trackIndex) { return begin(trackIndex) + counts[trackIndex]; }
125+
TrackExtensionCandidateN& get(int trackIndex, int candidateIndex) { return candidates[getFlatTrackExtensionCandidateIndex(trackIndex, candidateIndex)]; }
126+
const TrackExtensionCandidateN& get(int trackIndex, int candidateIndex) const { return candidates[getFlatTrackExtensionCandidateIndex(trackIndex, candidateIndex)]; }
127+
TrackExtensionCandidateN& getFlat(int flatIndex) { return candidates[flatIndex]; }
128+
129+
std::vector<TrackExtensionCandidateN> candidates;
130+
std::vector<int> counts;
131+
};
98132

99133
struct TrackFollowerScratch {
100134
std::vector<TrackExtensionHypothesis<NLayers>> activeHypotheses;

0 commit comments

Comments
 (0)