Skip to content

Commit 478e887

Browse files
mpucciompuccio
committed
ITS: allow CA seeding on configurable layer subsets
Add a SeedingLayers mask to TrackingParameters and use it when building the CA tracking topology. Transitions and cells are built only from layers that participate in seeding, while skipped non-seeding layers do not contribute to the effective seed length or hole budget. Move the track-seed selection predicate to TrackHelpers so CPU and GPU road finding use the same host/device logic.
1 parent 323df58 commit 478e887

9 files changed

Lines changed: 75 additions & 32 deletions

File tree

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -195,6 +195,7 @@ void processNeighboursHandler(const int startLevel,
195195
const int maxHoles,
196196
const int minTrackLength,
197197
const LayerMask holeLayerMask,
198+
const LayerMask nonCountingLayerMask,
198199
const std::vector<float>& layerxX0Host,
199200
const o2::base::Propagator* propagator,
200201
const o2::base::PropagatorF::MatCorrType matCorrType,

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -308,6 +308,7 @@ 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].getSeedingLayerMask();
311312
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
312313
bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
313314
for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) {
@@ -333,6 +334,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
333334
this->mTrkParams[iteration].MaxHoles,
334335
this->mTrkParams[iteration].MinTrackLength,
335336
this->mTrkParams[iteration].HoleLayerMask,
337+
nonCountingLayers,
336338
this->mTrkParams[iteration].LayerxX0,
337339
mTimeFrameGPU->getDevicePropagator(),
338340
this->mTrkParams[iteration].CorrType,

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

Lines changed: 5 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -32,8 +32,8 @@
3232
#include "ITStracking/Tracklet.h"
3333
#include "ITStracking/Cluster.h"
3434
#include "ITStracking/Cell.h"
35-
#include "ITStracking/TrackFollower.h"
3635
#include "ITStracking/TrackHelpers.h"
36+
#include "ITStracking/TrackFollower.h"
3737
#include "DataFormatsITS/TrackITS.h"
3838
#include "ITStrackingGPU/TrackingKernels.h"
3939
#include "ITStrackingGPU/Utils.h"
@@ -86,23 +86,6 @@ struct is_valid_pair {
8686
}
8787
};
8888

89-
template <int NLayers>
90-
struct seed_selector {
91-
float mMaxQ2Pt;
92-
float mMaxChi2;
93-
int mMaxHoles;
94-
int mMinTrackLength;
95-
LayerMask mHoleLayerMask;
96-
97-
GPUhd() seed_selector(float maxQ2Pt, float maxChi2, int maxHoles, int minTrackLength, LayerMask holeLayerMask) : mMaxQ2Pt(maxQ2Pt), mMaxChi2(maxChi2), mMaxHoles(maxHoles), mMinTrackLength(minTrackLength), mHoleLayerMask(holeLayerMask) {}
98-
GPUhd() bool operator()(const TrackSeed<NLayers>& seed) const
99-
{
100-
return !(seed.getQ2Pt() > mMaxQ2Pt || seed.getChi2() > mMaxChi2) &&
101-
seed.getHitLayerMask().length() >= mMinTrackLength &&
102-
seed.getHitLayerMask().isAllowed(mMaxHoles, mHoleLayerMask);
103-
}
104-
};
105-
10689
struct compare_track_chi2 {
10790
GPUhd() bool operator()(const TrackITSExt& a, const TrackITSExt& b) const
10891
{
@@ -977,6 +960,7 @@ void processNeighboursHandler(const int startLevel,
977960
const int maxHoles,
978961
const int minTrackLength,
979962
const LayerMask holeLayerMask,
963+
const LayerMask nonCountingLayerMask,
980964
const std::vector<float>& layerxX0Host,
981965
const o2::base::Propagator* propagator,
982966
const o2::base::PropagatorF::MatCorrType matCorrType,
@@ -1110,7 +1094,7 @@ void processNeighboursHandler(const int startLevel,
11101094
}
11111095
GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream));
11121096
thrust::device_vector<TrackSeed<NLayers>, gpu::TypedAllocator<TrackSeed<NLayers>>> outSeeds(updatedCellSeed.size(), allocTrackSeed);
1113-
auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), gpu::seed_selector<NLayers>(1.e3, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minTrackLength, holeLayerMask));
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});
11141098
auto s{end - outSeeds.begin()};
11151099
seedsHost.reserve(seedsHost.size() + s);
11161100
thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost));
@@ -1382,6 +1366,7 @@ template void processNeighboursHandler<7>(const int startLevel,
13821366
const int maxHoles,
13831367
const int minTrackLength,
13841368
const LayerMask holeLayerMask,
1369+
const LayerMask nonCountingLayerMask,
13851370
const std::vector<float>& layerxX0Host,
13861371
const o2::base::Propagator* propagator,
13871372
const o2::base::PropagatorF::MatCorrType matCorrType,
@@ -1582,6 +1567,7 @@ template void processNeighboursHandler<11>(const int startLevel,
15821567
const int maxHoles,
15831568
const int minTrackLength,
15841569
const LayerMask holeLayerMask,
1570+
const LayerMask nonCountingLayerMask,
15851571
const std::vector<float>& layerxX0Host,
15861572
const o2::base::Propagator* propagator,
15871573
const o2::base::PropagatorF::MatCorrType matCorrType,

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

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

4848
struct TrackingParameters {
49+
LayerMask getSeedingLayerMask() const noexcept
50+
{
51+
const auto layerSpan = LayerMask::span(0, NLayers - 1);
52+
return SeedingLayers.empty() ? layerSpan : (SeedingLayers & layerSpan);
53+
}
54+
55+
int getNSeedingLayers() const noexcept
56+
{
57+
return getSeedingLayerMask().count();
58+
}
59+
4960
int CellMinimumLevel() const noexcept
5061
{
5162
const int minClusters = MinTrackLength - (MaxHoles > 0 ? MaxHoles : 0);
5263
const int effectiveMinClusters = minClusters > constants::ClustersPerCell ? minClusters : constants::ClustersPerCell;
5364
return effectiveMinClusters - constants::ClustersPerCell + 1;
5465
}
55-
int NeighboursPerRoad() const noexcept { return NLayers - 3; }
56-
int CellsPerRoad() const noexcept { return NLayers - 2; }
57-
int TrackletsPerRoad() const noexcept { return NLayers - 1; }
66+
int NeighboursPerRoad() const noexcept { return getNSeedingLayers() - 3; }
67+
int CellsPerRoad() const noexcept { return getNSeedingLayers() - 2; }
68+
int TrackletsPerRoad() const noexcept { return getNSeedingLayers() - 1; }
5869
std::string asString() const;
5970

6071
IterationSteps PassFlags{IterationStep::FirstPass, IterationStep::RebuildClusterLUT};
@@ -76,6 +87,7 @@ struct TrackingParameters {
7687
int MinTrackLength = 7;
7788
int MaxHoles = 0;
7889
LayerMask HoleLayerMask = 0;
90+
LayerMask SeedingLayers = 0;
7991
float NSigmaCut = 5;
8092
float PVres = 1.e-2f;
8193
/// Trackleting cuts

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

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@
2121
#include "ITStracking/Cell.h"
2222
#include "ITStracking/Cluster.h"
2323
#include "ITStracking/Constants.h"
24+
#include "ITStracking/LayerMask.h"
2425
#include "ITStracking/MathUtils.h"
2526
#include "ITStracking/TrackITSInternal.h"
2627
#include "DetectorsBase/Propagator.h"
@@ -45,6 +46,37 @@ GPUhdi() bool isBetter(const o2::its::TrackITSInternal<NLayers>& a, const o2::it
4546
return isBetter(a.getNumberOfClusters(), a.getChi2(), b.getNumberOfClusters(), b.getChi2());
4647
}
4748

49+
GPUhdi() int getEffectiveTrackLength(LayerMask hitLayerMask, LayerMask nonCountingLayerMask)
50+
{
51+
if (hitLayerMask.empty()) {
52+
return 0;
53+
}
54+
return hitLayerMask.length() - (LayerMask::span(hitLayerMask.first(), hitLayerMask.last()) & nonCountingLayerMask).count();
55+
}
56+
57+
GPUhdi() LayerMask getEffectiveHoleMask(LayerMask hitLayerMask, LayerMask nonCountingLayerMask)
58+
{
59+
return hitLayerMask.holeMask() & ~nonCountingLayerMask;
60+
}
61+
62+
template <int NLayers>
63+
struct TrackSeedSelector {
64+
float maxQ2Pt;
65+
float maxChi2;
66+
int maxHoles;
67+
int minTrackLength;
68+
LayerMask holeLayerMask;
69+
LayerMask nonCountingLayerMask;
70+
71+
GPUhd() bool operator()(const TrackSeed<NLayers>& seed) const
72+
{
73+
const auto hitLayerMask = seed.getHitLayerMask();
74+
return !(seed.getQ2Pt() > maxQ2Pt || seed.getChi2() > maxChi2) &&
75+
getEffectiveTrackLength(hitLayerMask, nonCountingLayerMask) >= minTrackLength &&
76+
getEffectiveHoleMask(hitLayerMask, nonCountingLayerMask).isAllowedHoleMask(maxHoles, holeLayerMask);
77+
}
78+
};
79+
4880
// Find the populated interior layer closest to the radial midpoint.
4981
// If no layer can be found, return constants::UnusedIndex.
5082
// Should minimize the sagitta bias.

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

Lines changed: 13 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -104,15 +104,19 @@ class TrackingTopology
104104
#endif
105105
};
106106

107-
void init(int maxLayers, int maxHoles, Mask holeLayerMask)
107+
void init(int maxLayers, int maxHoles, Mask holeLayerMask, Mask seedingLayerMask = 0)
108108
{
109109
clear();
110110
mMaxLayers = o2::gpu::CAMath::Max(0, o2::gpu::CAMath::Min(maxLayers, NLayers));
111111
mMaxHoles = o2::gpu::CAMath::Max(maxHoles, 0);
112112
mHoleLayerMask = holeLayerMask;
113+
mSeedingLayerMask = seedingLayerMask.empty() ? Mask::span(0, mMaxLayers - 1) : (seedingLayerMask & Mask::span(0, mMaxLayers - 1));
113114
for (int fromLayer = 0; fromLayer < mMaxLayers; ++fromLayer) {
115+
if (!mSeedingLayerMask.has(fromLayer)) {
116+
continue;
117+
}
114118
for (int toLayer = fromLayer + 1; toLayer < mMaxLayers; ++toLayer) {
115-
if (Mask::skipped(fromLayer, toLayer).isAllowedHoleMask(mMaxHoles, mHoleLayerMask)) {
119+
if (mSeedingLayerMask.has(toLayer) && isAllowedSeedingLink(fromLayer, toLayer)) {
116120
mLinks[mNLinks++] = LayerLink{static_cast<Id>(fromLayer), static_cast<Id>(toLayer)};
117121
}
118122
}
@@ -126,7 +130,7 @@ class TrackingTopology
126130
continue;
127131
}
128132
const Mask hitMask{first.fromLayer, first.toLayer, second.toLayer};
129-
if (hitMask.isAllowed(mMaxHoles, mHoleLayerMask)) {
133+
if ((hitMask.holeMask() & mSeedingLayerMask).isAllowedHoleMask(mMaxHoles, mHoleLayerMask)) {
130134
mCells[mNCells++] = CellTopology{firstId, secondId, hitMask};
131135
}
132136
}
@@ -202,9 +206,15 @@ class TrackingTopology
202206
mNCellsByFirstLink = offset;
203207
}
204208

209+
bool isAllowedSeedingLink(int fromLayer, int toLayer) const noexcept
210+
{
211+
return (Mask::skipped(fromLayer, toLayer) & mSeedingLayerMask).isAllowedHoleMask(mMaxHoles, mHoleLayerMask);
212+
}
213+
205214
int mMaxLayers{0};
206215
int mMaxHoles{0};
207216
Mask mHoleLayerMask{0};
217+
Mask mSeedingLayerMask{0};
208218
Id mNLinks{0};
209219
Id mNCells{0};
210220
Id mNCellsByFirstLink{0};

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -59,6 +59,9 @@ std::string TrackingParameters::asString() const
5959
if (MaxHoles) {
6060
str += std::format(" MaxHoles:{} HoleMask:{}", MaxHoles, HoleLayerMask.asString());
6161
}
62+
if (!SeedingLayers.empty()) {
63+
str += std::format(" SeedingLayers:{}", SeedingLayers.asString());
64+
}
6265
if (PassFlags[IterationStep::TrackFollowerTop] || PassFlags[IterationStep::TrackFollowerBot]) {
6366
const bool top = PassFlags[IterationStep::TrackFollowerTop], bot = PassFlags[IterationStep::TrackFollowerBot];
6467
str += std::format(" TrackFollower:{} NSigmaZ/Phi:{:.2f}/{:.2f}",

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -249,7 +249,7 @@ 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);
252+
mDefaultTrackingTopology.init(maxLayers, trkParam.MaxHoles, trkParam.HoleLayerMask, trkParam.getSeedingLayerMask());
253253
}
254254

255255
template <int NLayers>
@@ -258,7 +258,7 @@ void TimeFrame<NLayers>::initTrackerTopologies(gsl::span<const TrackingParameter
258258
mTrackerTopologies.resize(trkParams.size());
259259
for (size_t iteration = 0; iteration < trkParams.size(); ++iteration) {
260260
const int iterationMaxLayers = std::min(maxLayers, trkParams[iteration].NLayers);
261-
mTrackerTopologies[iteration].init(iterationMaxLayers, trkParams[iteration].MaxHoles, trkParams[iteration].HoleLayerMask);
261+
mTrackerTopologies[iteration].init(iterationMaxLayers, trkParams[iteration].MaxHoles, trkParams[iteration].HoleLayerMask, trkParams[iteration].getSeedingLayerMask());
262262
}
263263
}
264264

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

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -759,13 +759,10 @@ 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].getSeedingLayerMask();
762763
for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
763764

764-
auto seedFilter = [&](const auto& seed) {
765-
return seed.getHitLayerMask().isAllowed(mTrkParams[iteration].MaxHoles, mTrkParams[iteration].HoleLayerMask) &&
766-
seed.getHitLayerMask().length() >= mTrkParams[iteration].MinTrackLength &&
767-
seed.getQ2Pt() <= 1.e3 && seed.getChi2() <= mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5);
768-
};
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};
769766

770767
bounded_vector<TrackSeedN> trackSeeds(mMemoryPool.get());
771768
for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) {

0 commit comments

Comments
 (0)