Skip to content

Commit 0528e00

Browse files
f3schf3sch
authored
ITS: fix rare valid lookup at the edge of acceptance (#15327)
* ITS: fix rare valid lookup at the edge of acceptance Technically, bin (0,0,0,0) is a valid result for the phi-z cluster query. Note though, that this if at all a super rare case and in local tests this did only showed up in the number of tracklets. So more a consistency fix than anything. Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: speedup line selection Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: improve logging Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: cleanup tracklet class Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: let staggering macro also run over non-staggered data Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: fix time assignment Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> --------- Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 7b7c9fa commit 0528e00

19 files changed

Lines changed: 313 additions & 222 deletions

‎Detectors/ITSMFT/ITS/macros/test/CheckStaggering.C‎

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -80,9 +80,9 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "")
8080
auto& ccdbmgr = o2::ccdb::BasicCCDBManager::instance();
8181
ccdbmgr.setURL("https://alice-ccdb.cern.ch");
8282
auto runDuration = ccdbmgr.getRunDuration(runNumber);
83-
auto tRun = runDuration.first + (runDuration.second - runDuration.first) / 2; // time stamp for the middle of the run duration
83+
auto tRun = runDuration.first + ((runDuration.second - runDuration.first) / 2); // time stamp for the middle of the run duration
8484
ccdbmgr.setTimestamp(tRun);
85-
printf("Run %d has TS %lld", runNumber, tRun);
85+
printf("Run %d has TS %ld", runNumber, tRun);
8686
auto geoAligned = ccdbmgr.get<TGeoManager>("GLO/Config/GeometryAligned");
8787
auto magField = ccdbmgr.get<o2::parameters::GRPMagField>("GLO/Config/GRPMagField");
8888
auto grpLHC = ccdbmgr.get<o2::parameters::GRPLHCIFData>("GLO/Config/GRPLHCIF");
@@ -134,7 +134,7 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "")
134134
auto hVtxZ = new TH1F("hVtxZ", "seeding vertices Z", 200, -16, 16);
135135
auto hVtxNCont = new TH1F("hVtxNCont", "seeding vertices contributors", 100, 0, 100);
136136
auto hVtxZNCont = new TProfile("hVtxZNCont", "seeding vertices z-contributors", 200, -16, 16);
137-
auto hVtxCls = new TProfile("hVtxCls", ";Cls/TF;Cls/Vtx", 400, 20000, 60000);
137+
auto hVtxCls = new TProfile("hVtxCls", ";Cls/TF;Cls/Vtx", 2000, 600000, 900000);
138138
auto hVtxTS = new TH1D("hVtxTS", "vtx time t0;t0 (BC)", o2::constants::lhc::LHCMaxBunches, 0, o2::constants::lhc::LHCMaxBunches);
139139

140140
const float minVtxWeight{5};
@@ -183,8 +183,12 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "")
183183

184184
tTrks->SetBranchAddress("ITSTrack", &trkArrPtr);
185185
tTrks->SetBranchAddress("Vertices", &vtxArrPtr);
186-
for (int i{0}; i < 7; ++i) {
187-
tCls->SetBranchAddress(Form("ITSClusterComp_%d", i), &clsArr[i]);
186+
if (tCls->GetBranchStatus("ITSClusterComp")) {
187+
tCls->SetBranchAddress("ITSClusterComp", &clsArr[0]);
188+
} else {
189+
for (int i{0}; i < 7; ++i) {
190+
tCls->SetBranchAddress(Form("ITSClusterComp_%d", i), &clsArr[i]);
191+
}
188192
}
189193

190194
for (int iTF{0}; tTrks->LoadTree(iTF) >= 0; ++iTF) {
@@ -193,7 +197,9 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "")
193197

194198
size_t ncls = 0;
195199
for (int i{0}; i < 7; ++i) {
196-
ncls += clsArr[i]->size();
200+
if (clsArr[i]) {
201+
ncls += clsArr[i]->size();
202+
}
197203
}
198204

199205
// for each TF built pool of positive and negaitve tracks

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

Lines changed: 3 additions & 17 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; }
@@ -408,7 +394,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel(
408394
const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + constants::Tolerance)}; /// protecting from overflows adding the detector resolution
409395
const float sigmaZ{o2::gpu::CAMath::Sqrt(math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInverseDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * MSAngle))};
410396
const int4 selectedBinsRect{o2::its::getBinsRect(currentCluster, layerIndex + 1, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut, *utils)};
411-
if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {
397+
if (selectedBinsRect.x < 0) {
412398
continue;
413399
}
414400
int phiBinsNum{selectedBinsRect.w - selectedBinsRect.y + 1};
@@ -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: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -56,18 +56,17 @@ 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;
6464
};
6565

6666
/// CellSeed: connections of three clusters
67-
class CellSeed final : public SeedBase<3>
67+
class CellSeed final : public SeedBase<constants::ClustersPerCell>
6868
{
69-
static constexpr int NStoredClusters = 3;
70-
using Base = SeedBase<NStoredClusters>;
69+
using Base = SeedBase<constants::ClustersPerCell>;
7170

7271
public:
7372
GPUhdDefault() CellSeed() = default;
@@ -98,7 +97,7 @@ class CellSeed final : public SeedBase<3>
9897
GPUhd() int getCluster(int layer) const
9998
{
10099
const int rel = layer - getInnerLayer();
101-
return (rel >= 0 && rel < NStoredClusters) ? this->clustersRaw()[rel] : constants::UnusedIndex;
100+
return (rel >= 0 && rel < constants::ClustersPerCell) ? this->clustersRaw()[rel] : constants::UnusedIndex;
102101
}
103102
};
104103

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

Lines changed: 8 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -19,9 +19,6 @@
1919
#include <array>
2020
#include <utility>
2121

22-
#include "GPUCommonDef.h"
23-
#include "GPUCommonDefAPI.h"
24-
2522
namespace o2::its::constants
2623
{
2724

@@ -30,12 +27,14 @@ constexpr float MB = KB * KB;
3027
constexpr float GB = MB * KB;
3128
constexpr bool DoTimeBenchmarks = true;
3229
constexpr bool SaveTimeBenchmarks = false;
33-
constexpr float Tolerance = 1e-12; // numerical tolerance
34-
constexpr int ClustersPerCell = 3;
35-
constexpr int UnusedIndex = -1;
36-
constexpr float Radl = 9.36f; // Radiation length of Si [cm]
37-
constexpr float Rho = 2.33f; // Density of Si [g/cm^3]
38-
constexpr int MaxIter = 4; // Max. supported iterations
30+
constexpr float Tolerance = 1e-12; // numerical tolerance
31+
constexpr int ClustersPerCell = 3; // number of clusters for a cell
32+
constexpr int UnusedIndex = -1; // global unused flag
33+
constexpr float UnsetValue = -999.f; // global unset value
34+
constexpr float Radl = 9.36f; // Radiation length of Si [cm]
35+
constexpr float Rho = 2.33f; // Density of Si [g/cm^3]
36+
constexpr int MaxIter = 4; // Max. supported iterations
37+
constexpr int MaxSelectedTrackletsPerCluster = 100; // vertexer: max lines per cluster
3938

4039
namespace helpers
4140
{

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

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,19 @@ struct LogLogThrottler {
5050
return false;
5151
}
5252
};
53+
54+
struct TimingStats {
55+
std::uint64_t calls = 0;
56+
double totalTimeMs = 0.;
57+
58+
void add(double timeMs)
59+
{
60+
++calls;
61+
totalTimeMs += timeMs;
62+
}
63+
double averageTimeMs() const { return calls ? totalTimeMs / static_cast<double>(calls) : 0.; }
64+
};
65+
5366
} // namespace o2::its
5467

55-
#endif
68+
#endif

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -124,7 +124,7 @@ GPUhdi() int4 getBinsRect(const Cluster& currentCluster, const int layerIndex,
124124

125125
if (zRangeMax < -utils.getLayerZ(layerIndex) ||
126126
zRangeMin > utils.getLayerZ(layerIndex) || zRangeMin > zRangeMax) {
127-
return int4{0, 0, 0, 0};
127+
return int4{-1, -1, -1, -1};
128128
}
129129

130130
return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex, zRangeMin)),

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -89,6 +89,12 @@ GPUhdi() float smallestAngleDifference(float a, float b)
8989
return o2::gpu::CAMath::Remainderf(b - a, o2::constants::math::TwoPI);
9090
}
9191

92+
GPUhdi() bool isPhiDifferenceBelow(const float phiA, const float phiB, const float phiCut)
93+
{
94+
const float deltaPhi = o2::gpu::CAMath::Abs(phiA - phiB);
95+
return deltaPhi < phiCut || deltaPhi > o2::constants::math::TwoPI - phiCut;
96+
}
97+
9298
GPUhdi() constexpr float Sq(float v)
9399
{
94100
return v * v;

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

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -26,9 +26,7 @@
2626
#include "ITStracking/Cell.h"
2727
#include "ITStracking/Cluster.h"
2828
#include "ITStracking/Configuration.h"
29-
#include "ITStracking/Constants.h"
3029
#include "ITStracking/ClusterLines.h"
31-
#include "ITStracking/Definitions.h"
3230
#include "ITStracking/Tracklet.h"
3331
#include "ITStracking/IndexTableUtils.h"
3432
#include "ITStracking/ExternalAllocator.h"
@@ -103,7 +101,7 @@ struct TimeFrame {
103101
void setBeamPosition(const float x, const float y, const float s2, const float base = 50.f, const float systematic = 0.f)
104102
{
105103
isBeamPositionOverridden = true;
106-
resetBeamXY(x, y, s2 / o2::gpu::CAMath::Sqrt(base * base + systematic));
104+
resetBeamXY(x, y, s2 / o2::gpu::CAMath::Sqrt((base * base) + systematic));
107105
}
108106

109107
float getBeamX() const { return mBeamPos[0]; }
@@ -249,7 +247,7 @@ struct TimeFrame {
249247

250248
// Propagator
251249
const o2::base::PropagatorImpl<float>* getDevicePropagator() const { return mPropagatorDevice; }
252-
virtual void setDevicePropagator(const o2::base::PropagatorImpl<float>*) {};
250+
virtual void setDevicePropagator(const o2::base::PropagatorImpl<float>* /*unused*/) {};
253251

254252
template <typename... T>
255253
void addClusterToLayer(int layer, T&&... args);

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

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -25,10 +25,12 @@
2525
#include <string_view>
2626
#include <utility>
2727
#include <sstream>
28+
#include <vector>
2829

2930
#include <oneapi/tbb/task_arena.h>
3031

3132
#include "ITStracking/Configuration.h"
33+
#include "ITStracking/Definitions.h"
3234
#include "ITStracking/TimeFrame.h"
3335
#include "ITStracking/TrackerTraits.h"
3436
#include "ITStracking/BoundedAllocator.h"
@@ -91,16 +93,18 @@ class Tracker
9193
double mTotalTime{0};
9294
std::shared_ptr<BoundedMemoryResource> mMemoryPool;
9395

94-
enum State {
96+
enum Steps {
9597
TFInit = 0,
9698
Trackleting,
9799
Celling,
98100
Neighbouring,
99101
Roading,
100-
NStates,
102+
NSteps,
101103
};
102-
State mCurState{TFInit};
103-
static constexpr std::array<const char*, NStates> StateNames{"TimeFrame initialisation", "Tracklet finding", "Cell finding", "Neighbour finding", "Road finding"};
104+
Steps mCurStep{TFInit};
105+
static constexpr std::array<const char*, NSteps> StateNames{"TimeFrame initialisation", "Tracklet finding", "Cell finding", "Neighbour finding", "Road finding"};
106+
std::vector<std::array<TimingStats, NSteps>> mTimingStats;
107+
void addTimingStatCurStep(int iteration, double timeMs);
104108
};
105109

106110
template <int NLayers>
@@ -125,22 +129,23 @@ float Tracker<NLayers>::evaluateTask(void (Tracker<NLayers>::*task)(T...), std::
125129
}
126130
logger(sstream.str());
127131

128-
if (mTrkParams[0].SaveTimeBenchmarks) {
132+
if (mTrkParams[iteration].SaveTimeBenchmarks) {
129133
std::string taskNameStr(taskName);
130134
std::transform(taskNameStr.begin(), taskNameStr.end(), taskNameStr.begin(),
131135
[](unsigned char c) { return std::tolower(c); });
132136
std::replace(taskNameStr.begin(), taskNameStr.end(), ' ', '_');
133137
if (std::ofstream file{"its_time_benchmarks.txt", std::ios::app}) {
134138
file << "trk:" << iteration << '\t' << taskNameStr << '\t' << diff << '\n';
135139
}
140+
addTimingStatCurStep(iteration, diff);
136141
}
137142

138143
} else {
139144
(this->*task)(std::forward<F>(args)...);
140145
}
141146

142147
if (mTrkParams[iteration].PrintMemory) {
143-
LOGP(info, "iter:{}:{}: {}", iteration, StateNames[mCurState], mMemoryPool->asString());
148+
LOGP(info, "iter:{}:{}: {}", iteration, StateNames[mCurStep], mMemoryPool->asString());
144149
}
145150

146151
return diff;

‎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_ */

0 commit comments

Comments
 (0)