Skip to content

Commit b74614f

Browse files
mpucciompuccio
committed
ITS: implement Felix' comments
1 parent f87bce2 commit b74614f

7 files changed

Lines changed: 38 additions & 36 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
@@ -195,7 +195,7 @@ void processNeighboursHandler(const int startLevel,
195195
const int maxHoles,
196196
const int minSeedingClusters,
197197
const LayerMask holeLayerMask,
198-
const LayerMask nonCountingLayerMask,
198+
const LayerMask nonSeedingLayerMask,
199199
const std::vector<float>& layerxX0Host,
200200
const o2::base::Propagator* propagator,
201201
const o2::base::PropagatorF::MatCorrType matCorrType,

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -308,7 +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 notSeedingLayers = this->mTrkParams[iteration].getNotSeedingLayerMask();
311+
const auto nonSeedingLayerMask = this->mTrkParams[iteration].getNonSeedingLayerMask();
312312
const auto minSeedingClusters = this->mTrkParams[iteration].getMinSeedingClusters();
313313
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
314314
bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
@@ -335,7 +335,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
335335
this->mTrkParams[iteration].MaxHoles,
336336
minSeedingClusters,
337337
this->mTrkParams[iteration].HoleLayerMask,
338-
notSeedingLayers,
338+
nonSeedingLayerMask,
339339
this->mTrkParams[iteration].LayerxX0,
340340
mTimeFrameGPU->getDevicePropagator(),
341341
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
@@ -960,7 +960,7 @@ void processNeighboursHandler(const int startLevel,
960960
const int maxHoles,
961961
const int minSeedingClusters,
962962
const LayerMask holeLayerMask,
963-
const LayerMask nonCountingLayerMask,
963+
const LayerMask nonSeedingLayerMask,
964964
const std::vector<float>& layerxX0Host,
965965
const o2::base::Propagator* propagator,
966966
const o2::base::PropagatorF::MatCorrType matCorrType,
@@ -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, minSeedingClusters, holeLayerMask, nonCountingLayerMask});
1097+
auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector<NLayers>{1.e3f, maxChi2NDF, startLevel, maxHoles, minSeedingClusters, holeLayerMask, nonSeedingLayerMask});
10981098
auto s{end - outSeeds.begin()};
10991099
seedsHost.reserve(seedsHost.size() + s);
11001100
thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost));
@@ -1366,7 +1366,7 @@ template void processNeighboursHandler<7>(const int startLevel,
13661366
const int maxHoles,
13671367
const int minSeedingClusters,
13681368
const LayerMask holeLayerMask,
1369-
const LayerMask nonCountingLayerMask,
1369+
const LayerMask nonSeedingLayerMask,
13701370
const std::vector<float>& layerxX0Host,
13711371
const o2::base::Propagator* propagator,
13721372
const o2::base::PropagatorF::MatCorrType matCorrType,
@@ -1567,7 +1567,7 @@ template void processNeighboursHandler<11>(const int startLevel,
15671567
const int maxHoles,
15681568
const int minSeedingClusters,
15691569
const LayerMask holeLayerMask,
1570-
const LayerMask nonCountingLayerMask,
1570+
const LayerMask nonSeedingLayerMask,
15711571
const std::vector<float>& layerxX0Host,
15721572
const o2::base::Propagator* propagator,
15731573
const o2::base::PropagatorF::MatCorrType matCorrType,

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -57,9 +57,9 @@ struct TrackingParameters {
5757
return SeedingLayers.empty() ? activeLayers : (SeedingLayers & activeLayers);
5858
}
5959

60-
LayerMask getNotSeedingLayerMask() const noexcept
60+
LayerMask getNonSeedingLayerMask() const noexcept
6161
{
62-
return ~getSeedingLayerMask();
62+
return ~(getSeedingLayerMask());
6363
}
6464

6565
int getNSeedingLayers() const noexcept

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

Lines changed: 22 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -35,45 +35,44 @@ GPUhdi() bool isBetter(const int nClustersA, const float chi2A, const int nClust
3535
return (nClustersA > nClustersB) || (nClustersA == nClustersB && chi2A < chi2B);
3636
}
3737

38-
GPUhdi() bool isBetter(const o2::its::TrackITS& a, const o2::its::TrackITS& b)
38+
GPUhdi() bool isBetter(const auto& a, const auto& b)
3939
{
4040
return isBetter(a.getNumberOfClusters(), a.getChi2(), b.getNumberOfClusters(), b.getChi2());
4141
}
4242

43-
template <int NLayers>
44-
GPUhdi() bool isBetter(const o2::its::TrackITSInternal<NLayers>& a, const o2::its::TrackITSInternal<NLayers>& b)
45-
{
46-
return isBetter(a.getNumberOfClusters(), a.getChi2(), b.getNumberOfClusters(), b.getChi2());
47-
}
48-
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-
6243
template <int NLayers>
6344
struct TrackSeedSelector {
6445
float maxQ2Pt;
6546
float maxChi2;
6647
int maxHoles;
6748
int minTrackLength;
6849
LayerMask holeLayerMask;
69-
LayerMask nonCountingLayerMask;
50+
LayerMask nonSeedingLayerMask;
51+
52+
GPUhd() TrackSeedSelector(float maxQ2Pt, float maxChi2NDF, int startLevel, int maxHoles, int minTrackLength, LayerMask holeLayerMask, LayerMask nonSeedingLayerMask)
53+
: maxQ2Pt{maxQ2Pt}, maxChi2{maxChi2NDF * ((startLevel + 2) * 2 - 5)}, maxHoles{maxHoles}, minTrackLength{minTrackLength}, holeLayerMask{holeLayerMask}, nonSeedingLayerMask{nonSeedingLayerMask}
54+
{
55+
}
56+
57+
static GPUhdi() int getEffectiveTrackLength(LayerMask hitLayerMask, LayerMask excludedLayerMask)
58+
{
59+
if (hitLayerMask.empty()) {
60+
return 0;
61+
}
62+
return hitLayerMask.length() - (LayerMask::span(hitLayerMask.first(), hitLayerMask.last()) & excludedLayerMask).count();
63+
}
64+
65+
static GPUhdi() LayerMask getEffectiveHoleMask(LayerMask hitLayerMask, LayerMask excludedLayerMask)
66+
{
67+
return hitLayerMask.holeMask() & ~excludedLayerMask;
68+
}
7069

7170
GPUhd() bool operator()(const TrackSeed<NLayers>& seed) const
7271
{
7372
const auto hitLayerMask = seed.getHitLayerMask();
7473
return !(seed.getQ2Pt() > maxQ2Pt || seed.getChi2() > maxChi2) &&
75-
getEffectiveTrackLength(hitLayerMask, nonCountingLayerMask) >= minTrackLength &&
76-
getEffectiveHoleMask(hitLayerMask, nonCountingLayerMask).isAllowedHoleMask(maxHoles, holeLayerMask);
74+
getEffectiveTrackLength(hitLayerMask, nonSeedingLayerMask) >= minTrackLength &&
75+
getEffectiveHoleMask(hitLayerMask, nonSeedingLayerMask).isAllowedHoleMask(maxHoles, holeLayerMask);
7776
}
7877
};
7978

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

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -70,6 +70,7 @@ class TrackingTopology
7070
const CellTopology* cells{nullptr};
7171
const Range* cellsByFirstLinkIndex{nullptr};
7272
const Id* cellsByFirstLink{nullptr};
73+
Mask seedingLayerMask{0};
7374
Id nLinks{0};
7475
Id nCells{0};
7576
Id nCellsByFirstLink{0};
@@ -81,7 +82,7 @@ class TrackingTopology
8182
#ifndef GPUCA_GPUCODE
8283
std::string asString() const
8384
{
84-
std::string out = fmt::format("TrackingTopology: links={} cells={}", nLinks, nCells);
85+
std::string out = fmt::format("TrackingTopology: links={} cells={} seedingLayers={}", nLinks, nCells, seedingLayerMask.asString());
8586
out += "\n links:";
8687
for (Id linkId = 0; linkId < nLinks; ++linkId) {
8788
const auto& t = links[linkId];
@@ -150,6 +151,7 @@ class TrackingTopology
150151
mCells.data(),
151152
mCellsByFirstLinkIndex.data(),
152153
mCellsByFirstLink.data(),
154+
mSeedingLayerMask,
153155
mNLinks,
154156
mNCells,
155157
mNCellsByFirstLink};
@@ -164,6 +166,7 @@ class TrackingTopology
164166
deviceCells,
165167
deviceCellsByFirstLinkIndex,
166168
deviceCellsByFirstLink,
169+
mSeedingLayerMask,
167170
mNLinks,
168171
mNCells,
169172
mNCellsByFirstLink};

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -759,11 +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 notSeedingLayers = mTrkParams[iteration].getNotSeedingLayerMask();
762+
const auto nonSeedingLayerMask = mTrkParams[iteration].getNonSeedingLayerMask();
763763
const auto minSeedingClusters = mTrkParams[iteration].getMinSeedingClusters();
764764
for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
765765

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

768768
bounded_vector<TrackSeedN> trackSeeds(mMemoryPool.get());
769769
for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) {
@@ -875,7 +875,7 @@ void TrackerTraits<NLayers>::acceptTracks(int iteration,
875875
}
876876

877877
/// 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) {
878+
if (track::TrackSeedSelector<NLayers>::getEffectiveTrackLength(hitLayerMask, inactiveLayerMask) < minTrackLength) {
879879
continue;
880880
}
881881

0 commit comments

Comments
 (0)