Skip to content

Commit 5bd08b4

Browse files
f3schf3sch
committed
ITS: cleanup tracklet class
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 3953e79 commit 5bd08b4

6 files changed

Lines changed: 41 additions & 87 deletions

File tree

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

Lines changed: 2 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -47,20 +47,6 @@ namespace o2::its
4747
namespace gpu
4848
{
4949

50-
struct sort_tracklets {
51-
GPUhd() bool operator()(const Tracklet& a, const Tracklet& b)
52-
{
53-
if (a.firstClusterIndex != b.firstClusterIndex) {
54-
return a.firstClusterIndex < b.firstClusterIndex;
55-
}
56-
return a.secondClusterIndex < b.secondClusterIndex;
57-
}
58-
};
59-
60-
struct equal_tracklets {
61-
GPUhd() bool operator()(const Tracklet& a, const Tracklet& b) { return a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex == b.secondClusterIndex; }
62-
};
63-
6450
template <typename T1, typename T2>
6551
struct sort_by_second {
6652
GPUhd() bool operator()(const gpuPair<T1, T2>& a, const gpuPair<T1, T2>& b) const { return a.second < b.second; }
@@ -687,8 +673,8 @@ void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
687673
mulScatAng[layer]);
688674
thrust::device_ptr<Tracklet> tracklets_ptr(spanTracklets[layer]);
689675
auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator<char>(alloc)).on(streams[layer].get());
690-
thrust::sort(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer], gpu::sort_tracklets());
691-
auto unique_end = thrust::unique(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer], gpu::equal_tracklets());
676+
thrust::sort(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer]);
677+
auto unique_end = thrust::unique(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer]);
692678
nTracklets[layer] = unique_end - tracklets_ptr;
693679
if (layer) {
694680
GPUChkErrS(cudaMemsetAsync(trackletsLUTsHost[layer], 0, (nClusters[layer] + 1) * sizeof(int), streams[layer].get()));

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -56,8 +56,8 @@ class SeedBase : public o2::track::TrackParCovF
5656
GPUhd() const auto& clustersRaw() const { return mClusters; }
5757

5858
private:
59-
float mChi2 = -999.f;
60-
int mLevel = constants::UnusedIndex;
59+
float mChi2{constants::UnsetValue};
60+
int mLevel{constants::UnusedIndex};
6161
std::array<int, 2> mTracklets = constants::helpers::initArray<int, 2, constants::UnusedIndex>();
6262
std::array<int, NClusters> mClusters = constants::helpers::initArray<int, NClusters, constants::UnusedIndex>();
6363
TimeEstBC mTime;

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,7 @@ constexpr bool SaveTimeBenchmarks = false;
3030
constexpr float Tolerance = 1e-12; // numerical tolerance
3131
constexpr int ClustersPerCell = 3; // number of clusters for a cell
3232
constexpr int UnusedIndex = -1; // global unused flag
33+
constexpr float UnsetValue = -999.f; // global unset value
3334
constexpr float Radl = 9.36f; // Radiation length of Si [cm]
3435
constexpr float Rho = 2.33f; // Density of Si [g/cm^3]
3536
constexpr int MaxIter = 4; // Max. supported iterations

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

Lines changed: 23 additions & 46 deletions
Original file line numberDiff line numberDiff line change
@@ -24,27 +24,35 @@
2424
#include "GPUCommonDef.h"
2525
#include "GPUCommonLogger.h"
2626

27-
#ifndef GPUCA_GPUCODE_DEVICE
28-
#ifndef GPU_NO_FMT
29-
#include <string>
30-
#include <fmt/format.h>
31-
#endif
32-
#endif
33-
3427
namespace o2::its
3528
{
3629

30+
// tracklets are entirely determined by their two cluster idx
3731
struct Tracklet final {
3832
GPUhdDefault() Tracklet() = default;
39-
GPUhdi() Tracklet(const int, const int, const Cluster&, const Cluster&, const TimeEstBC& t);
40-
GPUhdi() Tracklet(const int, const int, float tanL, float phi, const TimeEstBC& t);
41-
GPUhdDefault() bool operator==(const Tracklet&) const = default;
42-
GPUhdi() unsigned char isEmpty() const
33+
GPUhdi() Tracklet(const int firstClusterOrderingIndex, const int secondClusterOrderingIndex,
34+
const Cluster& firstCluster, const Cluster& secondCluster, const TimeEstBC& t)
35+
: firstClusterIndex(firstClusterOrderingIndex),
36+
secondClusterIndex(secondClusterOrderingIndex),
37+
tanLambda((firstCluster.zCoordinate - secondCluster.zCoordinate) / (firstCluster.radius - secondCluster.radius)),
38+
phi(o2::gpu::GPUCommonMath::ATan2(firstCluster.yCoordinate - secondCluster.yCoordinate, firstCluster.xCoordinate - secondCluster.xCoordinate)),
39+
mTime(t) {}
40+
41+
GPUhdi() Tracklet(const int idx0, const int idx1, float tanL, float phi, const TimeEstBC& t)
42+
: firstClusterIndex(idx0),
43+
secondClusterIndex(idx1),
44+
tanLambda(tanL),
45+
phi(phi),
46+
mTime(t) {}
47+
GPUhdi() bool operator<(const Tracklet& o) const noexcept
48+
{
49+
return (firstClusterIndex != o.firstClusterIndex) ? firstClusterIndex < o.firstClusterIndex : secondClusterIndex < o.secondClusterIndex;
50+
}
51+
GPUhdi() bool operator==(const Tracklet& o) const noexcept
4352
{
44-
return firstClusterIndex < 0 || secondClusterIndex < 0;
53+
return firstClusterIndex == o.firstClusterIndex && secondClusterIndex == o.secondClusterIndex;
4554
}
4655
GPUhdi() bool isCompatible(const Tracklet& o) const { return mTime.isCompatible(o.mTime); }
47-
GPUhdi() unsigned char operator<(const Tracklet&) const;
4856
GPUhd() void print() const
4957
{
5058
LOGP(info, "TRKLT: fClIdx:{} sClIdx:{} ts:{}+/-{} TgL={} Phi={}", firstClusterIndex, secondClusterIndex, mTime.getTimeStamp(), mTime.getTimeStampError(), tanLambda, phi);
@@ -54,44 +62,13 @@ struct Tracklet final {
5462

5563
int firstClusterIndex{constants::UnusedIndex};
5664
int secondClusterIndex{constants::UnusedIndex};
57-
float tanLambda{-999};
58-
float phi{-999};
65+
float tanLambda{constants::UnsetValue};
66+
float phi{constants::UnsetValue};
5967
TimeEstBC mTime;
6068

6169
ClassDefNV(Tracklet, 1);
6270
};
6371

64-
GPUhdi() Tracklet::Tracklet(const int firstClusterOrderingIndex, const int secondClusterOrderingIndex,
65-
const Cluster& firstCluster, const Cluster& secondCluster, const TimeEstBC& t)
66-
: firstClusterIndex(firstClusterOrderingIndex),
67-
secondClusterIndex(secondClusterOrderingIndex),
68-
tanLambda((firstCluster.zCoordinate - secondCluster.zCoordinate) /
69-
(firstCluster.radius - secondCluster.radius)),
70-
phi(o2::gpu::GPUCommonMath::ATan2(firstCluster.yCoordinate - secondCluster.yCoordinate,
71-
firstCluster.xCoordinate - secondCluster.xCoordinate)),
72-
mTime(t)
73-
{
74-
// Nothing to do
75-
}
76-
77-
GPUhdi() Tracklet::Tracklet(const int idx0, const int idx1, float tanL, float phi, const TimeEstBC& t)
78-
: firstClusterIndex(idx0),
79-
secondClusterIndex(idx1),
80-
tanLambda(tanL),
81-
phi(phi),
82-
mTime(t)
83-
{
84-
// Nothing to do
85-
}
86-
87-
GPUhdi() unsigned char Tracklet::operator<(const Tracklet& t) const
88-
{
89-
if (isEmpty()) {
90-
return false;
91-
}
92-
return true;
93-
}
94-
9572
} // namespace o2::its
9673

9774
#endif /* TRACKINGITS_INCLUDE_TRACKLET_H_ */

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

Lines changed: 4 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -206,19 +206,11 @@ void TrackerTraits<NLayers>::computeLayerTracklets(const int iteration, int iVer
206206
}
207207

208208
tbb::parallel_for(0, mTrkParams[iteration].TrackletsPerRoad(), [&](const int iLayer) {
209-
/// Sort tracklets
209+
/// Sort tracklets & remove duplicates
210+
// duplicates can exist simply since we evaluate per vertex
210211
auto& trkl{mTimeFrame->getTracklets()[iLayer]};
211-
std::sort(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) -> bool {
212-
if (a.firstClusterIndex != b.firstClusterIndex) {
213-
return a.firstClusterIndex < b.firstClusterIndex;
214-
}
215-
return a.secondClusterIndex < b.secondClusterIndex;
216-
});
217-
/// Remove duplicates
218-
trkl.erase(std::unique(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) -> bool {
219-
return a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex == b.secondClusterIndex;
220-
}),
221-
trkl.end());
212+
std::sort(trkl.begin(), trkl.end());
213+
trkl.erase(std::unique(trkl.begin(), trkl.end()), trkl.end());
222214
trkl.shrink_to_fit();
223215
if (iLayer > 0) { /// recalculate lut
224216
auto& lut{mTimeFrame->getTrackletsLookupTable()[iLayer - 1]};

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

Lines changed: 9 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -269,20 +269,18 @@ void VertexerTraits<NLayers>::computeTracklets(const int iteration)
269269
if (mTimeFrame->hasMCinformation()) {
270270
for (const auto& trk : mTimeFrame->getTracklets()[0]) {
271271
o2::MCCompLabel label;
272-
if (!trk.isEmpty()) {
273-
int sortedId0{trk.firstClusterIndex};
274-
int sortedId1{trk.secondClusterIndex};
275-
for (const auto& lab0 : mTimeFrame->getClusterLabels(0, mTimeFrame->getClusters()[0][sortedId0].clusterId)) {
276-
for (const auto& lab1 : mTimeFrame->getClusterLabels(1, mTimeFrame->getClusters()[1][sortedId1].clusterId)) {
277-
if (lab0 == lab1 && lab0.isValid()) {
278-
label = lab0;
279-
break;
280-
}
281-
}
282-
if (label.isValid()) {
272+
int sortedId0{trk.firstClusterIndex};
273+
int sortedId1{trk.secondClusterIndex};
274+
for (const auto& lab0 : mTimeFrame->getClusterLabels(0, mTimeFrame->getClusters()[0][sortedId0].clusterId)) {
275+
for (const auto& lab1 : mTimeFrame->getClusterLabels(1, mTimeFrame->getClusters()[1][sortedId1].clusterId)) {
276+
if (lab0 == lab1 && lab0.isValid()) {
277+
label = lab0;
283278
break;
284279
}
285280
}
281+
if (label.isValid()) {
282+
break;
283+
}
286284
}
287285
mTimeFrame->getTrackletsLabel(0).emplace_back(label);
288286
}

0 commit comments

Comments
 (0)