Skip to content

Commit f87bce2

Browse files
mpucciompuccio
committed
ITS: make sure that seeds are not cut by min track lenght
1 parent fc218b9 commit f87bce2

7 files changed

Lines changed: 59 additions & 18 deletions

File tree

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -193,7 +193,7 @@ void processNeighboursHandler(const int startLevel,
193193
const float MaxChi2ClusterAttachment,
194194
const float maxChi2NDF,
195195
const int maxHoles,
196-
const int minTrackLength,
196+
const int minSeedingClusters,
197197
const LayerMask holeLayerMask,
198198
const LayerMask nonCountingLayerMask,
199199
const std::vector<float>& layerxX0Host,

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

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -308,7 +308,8 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
308308
const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop];
309309
const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot];
310310
const bool extendTracks = extendTop || extendBot;
311-
const auto nonCountingLayers = this->mTrkParams[iteration].InactiveLayerMask | ~this->mTrkParams[iteration].getSeedingLayerMask();
311+
const auto notSeedingLayers = this->mTrkParams[iteration].getNotSeedingLayerMask();
312+
const auto minSeedingClusters = this->mTrkParams[iteration].getMinSeedingClusters();
312313
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
313314
bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
314315
for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) {
@@ -332,9 +333,9 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
332333
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
333334
this->mTrkParams[iteration].MaxChi2NDF,
334335
this->mTrkParams[iteration].MaxHoles,
335-
this->mTrkParams[iteration].MinTrackLength,
336+
minSeedingClusters,
336337
this->mTrkParams[iteration].HoleLayerMask,
337-
nonCountingLayers,
338+
notSeedingLayers,
338339
this->mTrkParams[iteration].LayerxX0,
339340
mTimeFrameGPU->getDevicePropagator(),
340341
this->mTrkParams[iteration].CorrType,

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -958,7 +958,7 @@ void processNeighboursHandler(const int startLevel,
958958
const float maxChi2ClusterAttachment,
959959
const float maxChi2NDF,
960960
const int maxHoles,
961-
const int minTrackLength,
961+
const int minSeedingClusters,
962962
const LayerMask holeLayerMask,
963963
const LayerMask nonCountingLayerMask,
964964
const std::vector<float>& layerxX0Host,
@@ -1094,7 +1094,7 @@ void processNeighboursHandler(const int startLevel,
10941094
}
10951095
GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream));
10961096
thrust::device_vector<TrackSeed<NLayers>, gpu::TypedAllocator<TrackSeed<NLayers>>> outSeeds(updatedCellSeed.size(), allocTrackSeed);
1097-
auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector<NLayers>{1.e3f, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minTrackLength, holeLayerMask, nonCountingLayerMask});
1097+
auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector<NLayers>{1.e3f, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minSeedingClusters, holeLayerMask, nonCountingLayerMask});
10981098
auto s{end - outSeeds.begin()};
10991099
seedsHost.reserve(seedsHost.size() + s);
11001100
thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost));
@@ -1364,7 +1364,7 @@ template void processNeighboursHandler<7>(const int startLevel,
13641364
const float maxChi2ClusterAttachment,
13651365
const float maxChi2NDF,
13661366
const int maxHoles,
1367-
const int minTrackLength,
1367+
const int minSeedingClusters,
13681368
const LayerMask holeLayerMask,
13691369
const LayerMask nonCountingLayerMask,
13701370
const std::vector<float>& layerxX0Host,
@@ -1565,7 +1565,7 @@ template void processNeighboursHandler<11>(const int startLevel,
15651565
const float maxChi2ClusterAttachment,
15661566
const float maxChi2NDF,
15671567
const int maxHoles,
1568-
const int minTrackLength,
1568+
const int minSeedingClusters,
15691569
const LayerMask holeLayerMask,
15701570
const LayerMask nonCountingLayerMask,
15711571
const std::vector<float>& layerxX0Host,

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

Lines changed: 21 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -46,22 +46,38 @@ enum class IterationStep : uint16_t {
4646
using IterationSteps = o2::utils::EnumFlags<IterationStep>;
4747

4848
struct TrackingParameters {
49+
LayerMask getActiveLayerMask() const noexcept
50+
{
51+
return LayerMask::span(0, NLayers - 1) & ~InactiveLayerMask;
52+
}
53+
4954
LayerMask getSeedingLayerMask() const noexcept
5055
{
51-
const auto layerSpan = LayerMask::span(0, NLayers - 1);
52-
return SeedingLayers.empty() ? layerSpan : (SeedingLayers & layerSpan);
56+
const auto activeLayers = getActiveLayerMask();
57+
return SeedingLayers.empty() ? activeLayers : (SeedingLayers & activeLayers);
58+
}
59+
60+
LayerMask getNotSeedingLayerMask() const noexcept
61+
{
62+
return ~getSeedingLayerMask();
5363
}
5464

5565
int getNSeedingLayers() const noexcept
5666
{
5767
return getSeedingLayerMask().count();
5868
}
5969

60-
int CellMinimumLevel() const noexcept
70+
int getMinSeedingClusters() const noexcept
6171
{
6272
const int minClusters = MinTrackLength - (MaxHoles > 0 ? MaxHoles : 0);
63-
const int effectiveMinClusters = minClusters > constants::ClustersPerCell ? minClusters : constants::ClustersPerCell;
64-
return effectiveMinClusters - constants::ClustersPerCell + 1;
73+
const int minClustersWithCells = minClusters > constants::ClustersPerCell ? minClusters : constants::ClustersPerCell;
74+
const int nSeedingLayers = getNSeedingLayers();
75+
return minClustersWithCells < nSeedingLayers ? minClustersWithCells : nSeedingLayers;
76+
}
77+
78+
int CellMinimumLevel() const noexcept
79+
{
80+
return getMinSeedingClusters() - constants::ClustersPerCell + 1;
6581
}
6682
int NeighboursPerRoad() const noexcept { return getNSeedingLayers() - 3; }
6783
int CellsPerRoad() const noexcept { return getNSeedingLayers() - 2; }

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -111,6 +111,11 @@ class TrackingTopology
111111
mMaxHoles = o2::gpu::CAMath::Max(maxHoles, 0);
112112
mHoleLayerMask = holeLayerMask;
113113
mSeedingLayerMask = seedingLayerMask.empty() ? Mask::span(0, mMaxLayers - 1) : (seedingLayerMask & Mask::span(0, mMaxLayers - 1));
114+
#ifndef GPUCA_GPUCODE
115+
if (mSeedingLayerMask.count() < constants::ClustersPerCell) {
116+
LOGP(fatal, "Tracking topology has {} seeding layers, but at least {} are required to build CA cells", mSeedingLayerMask.count(), constants::ClustersPerCell);
117+
}
118+
#endif
114119
for (int fromLayer = 0; fromLayer < mMaxLayers; ++fromLayer) {
115120
if (!mSeedingLayerMask.has(fromLayer)) {
116121
continue;

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

Lines changed: 12 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -249,16 +249,25 @@ void TimeFrame<NLayers>::initVertexingTopology(const TrackingParameters& trkPara
249249
template <int NLayers>
250250
void TimeFrame<NLayers>::initDefaultTrackingTopology(const TrackingParameters& trkParam, const int maxLayers)
251251
{
252-
mDefaultTrackingTopology.init(maxLayers, trkParam.MaxHoles, trkParam.HoleLayerMask, trkParam.getSeedingLayerMask());
252+
if (maxLayers < trkParam.NLayers) {
253+
LOGP(fatal, "Default tracking topology limited to {} layers, but the tracking parameters expect {}", maxLayers, trkParam.NLayers);
254+
}
255+
mDefaultTrackingTopology.init(trkParam.NLayers, trkParam.MaxHoles, trkParam.HoleLayerMask, trkParam.getSeedingLayerMask());
253256
}
254257

255258
template <int NLayers>
256259
void TimeFrame<NLayers>::initTrackerTopologies(gsl::span<const TrackingParameters> trkParams, const int maxLayers)
257260
{
258261
mTrackerTopologies.resize(trkParams.size());
259262
for (size_t iteration = 0; iteration < trkParams.size(); ++iteration) {
260-
const int iterationMaxLayers = std::min(maxLayers, trkParams[iteration].NLayers);
261-
mTrackerTopologies[iteration].init(iterationMaxLayers, trkParams[iteration].MaxHoles, trkParams[iteration].HoleLayerMask, trkParams[iteration].getSeedingLayerMask());
263+
if (maxLayers < trkParams[iteration].NLayers) {
264+
LOGP(fatal, "Iteration {}: tracking topology limited to {} layers, but the tracking parameters expect {}", iteration, maxLayers, trkParams[iteration].NLayers);
265+
}
266+
const int nActiveLayers = trkParams[iteration].getActiveLayerMask().count();
267+
if (trkParams[iteration].MinTrackLength > nActiveLayers) {
268+
LOGP(fatal, "Iteration {}: MinTrackLength {} cannot be satisfied with {} active layers", iteration, trkParams[iteration].MinTrackLength, nActiveLayers);
269+
}
270+
mTrackerTopologies[iteration].init(trkParams[iteration].NLayers, trkParams[iteration].MaxHoles, trkParams[iteration].HoleLayerMask, trkParams[iteration].getSeedingLayerMask());
262271
}
263272
}
264273

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

Lines changed: 12 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -759,10 +759,11 @@ void TrackerTraits<NLayers>::findRoads(const int iteration)
759759
unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data();
760760
}
761761
const auto topology = mTimeFrame->getTrackingTopologyView();
762-
const auto nonCountingLayers = mTrkParams[iteration].InactiveLayerMask | ~mTrkParams[iteration].getSeedingLayerMask();
762+
const auto notSeedingLayers = mTrkParams[iteration].getNotSeedingLayerMask();
763+
const auto minSeedingClusters = mTrkParams[iteration].getMinSeedingClusters();
763764
for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
764765

765-
const track::TrackSeedSelector<NLayers> seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, mTrkParams[iteration].MinTrackLength, mTrkParams[iteration].HoleLayerMask, nonCountingLayers};
766+
const track::TrackSeedSelector<NLayers> seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, minSeedingClusters, mTrkParams[iteration].HoleLayerMask, notSeedingLayers};
766767

767768
bounded_vector<TrackSeedN> trackSeeds(mMemoryPool.get());
768769
for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) {
@@ -852,14 +853,18 @@ void TrackerTraits<NLayers>::acceptTracks(int iteration,
852853
auto& trks = mTimeFrame->getTracks();
853854
trks.reserve(trks.size() + tracks.size());
854855
const float smallestROFHalf = mTimeFrame->getROFOverlapTableView().getClockLayer().mROFLength * 0.5f;
856+
const int minTrackLength = mTrkParams[iteration].MinTrackLength;
857+
const LayerMask inactiveLayerMask = mTrkParams[iteration].InactiveLayerMask;
855858
for (auto& track : tracks) {
856859
int nShared = 0;
857860
bool isFirstShared{false};
858861
int firstLayer{-1}, firstCluster{-1};
862+
LayerMask hitLayerMask{0};
859863
for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) {
860864
if (track.getClusterIndex(iLayer) == constants::UnusedIndex) {
861865
continue;
862866
}
867+
hitLayerMask.set(iLayer);
863868
bool isShared = mTimeFrame->isClusterUsed(iLayer, track.getClusterIndex(iLayer));
864869
nShared += int(isShared);
865870
if (firstLayer < 0) {
@@ -869,6 +874,11 @@ void TrackerTraits<NLayers>::acceptTracks(int iteration,
869874
}
870875
}
871876

877+
/// seeds are selected with a length cut relaxed to the seeding layers: enforce the full minimum length before accepting the final track
878+
if (track::getEffectiveTrackLength(hitLayerMask, inactiveLayerMask) < minTrackLength) {
879+
continue;
880+
}
881+
872882
/// do not account for the first cluster in the shared clusters number if it is allowed
873883
if (nShared - int(isFirstShared && mTrkParams[iteration].AllowSharingFirstCluster) > mTrkParams[iteration].SharedMaxClusters) {
874884
continue;

0 commit comments

Comments
 (0)