From 478e887c00e9aaaa6c3a98564a122cb21517bd6f Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Tue, 9 Jun 2026 20:02:26 +0200 Subject: [PATCH 1/5] 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. --- .../GPU/ITStrackingGPU/TrackingKernels.h | 1 + .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 2 ++ .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 24 +++----------- .../include/ITStracking/Configuration.h | 18 +++++++++-- .../include/ITStracking/TrackHelpers.h | 32 +++++++++++++++++++ .../include/ITStracking/TrackingTopology.h | 16 ++++++++-- .../ITSMFT/ITS/tracking/src/Configuration.cxx | 3 ++ .../ITSMFT/ITS/tracking/src/TimeFrame.cxx | 4 +-- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 7 ++-- 9 files changed, 75 insertions(+), 32 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 7ad15ec9ddd12..e0d53766f08d3 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -195,6 +195,7 @@ void processNeighboursHandler(const int startLevel, const int maxHoles, const int minTrackLength, const LayerMask holeLayerMask, + const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index a19759d0577ec..4d43b765ef325 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -308,6 +308,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; const bool extendTracks = extendTop || extendBot; + const auto nonCountingLayers = ~this->mTrkParams[iteration].getSeedingLayerMask(); for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { bounded_vector> trackSeeds(this->getMemoryPool().get()); for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) { @@ -333,6 +334,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].MaxHoles, this->mTrkParams[iteration].MinTrackLength, this->mTrkParams[iteration].HoleLayerMask, + nonCountingLayers, this->mTrkParams[iteration].LayerxX0, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[iteration].CorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index cdeb95c67be30..1f974bfa3cda2 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -32,8 +32,8 @@ #include "ITStracking/Tracklet.h" #include "ITStracking/Cluster.h" #include "ITStracking/Cell.h" -#include "ITStracking/TrackFollower.h" #include "ITStracking/TrackHelpers.h" +#include "ITStracking/TrackFollower.h" #include "DataFormatsITS/TrackITS.h" #include "ITStrackingGPU/TrackingKernels.h" #include "ITStrackingGPU/Utils.h" @@ -86,23 +86,6 @@ struct is_valid_pair { } }; -template -struct seed_selector { - float mMaxQ2Pt; - float mMaxChi2; - int mMaxHoles; - int mMinTrackLength; - LayerMask mHoleLayerMask; - - GPUhd() seed_selector(float maxQ2Pt, float maxChi2, int maxHoles, int minTrackLength, LayerMask holeLayerMask) : mMaxQ2Pt(maxQ2Pt), mMaxChi2(maxChi2), mMaxHoles(maxHoles), mMinTrackLength(minTrackLength), mHoleLayerMask(holeLayerMask) {} - GPUhd() bool operator()(const TrackSeed& seed) const - { - return !(seed.getQ2Pt() > mMaxQ2Pt || seed.getChi2() > mMaxChi2) && - seed.getHitLayerMask().length() >= mMinTrackLength && - seed.getHitLayerMask().isAllowed(mMaxHoles, mHoleLayerMask); - } -}; - struct compare_track_chi2 { GPUhd() bool operator()(const TrackITSExt& a, const TrackITSExt& b) const { @@ -977,6 +960,7 @@ void processNeighboursHandler(const int startLevel, const int maxHoles, const int minTrackLength, const LayerMask holeLayerMask, + const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, @@ -1110,7 +1094,7 @@ void processNeighboursHandler(const int startLevel, } GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream)); thrust::device_vector, gpu::TypedAllocator>> outSeeds(updatedCellSeed.size(), allocTrackSeed); - auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), gpu::seed_selector(1.e3, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minTrackLength, holeLayerMask)); + auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{1.e3f, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minTrackLength, holeLayerMask, nonCountingLayerMask}); auto s{end - outSeeds.begin()}; seedsHost.reserve(seedsHost.size() + s); thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost)); @@ -1382,6 +1366,7 @@ template void processNeighboursHandler<7>(const int startLevel, const int maxHoles, const int minTrackLength, const LayerMask holeLayerMask, + const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, @@ -1582,6 +1567,7 @@ template void processNeighboursHandler<11>(const int startLevel, const int maxHoles, const int minTrackLength, const LayerMask holeLayerMask, + const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index ee8baf20e66c9..61ca263230bb9 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -46,15 +46,26 @@ enum class IterationStep : uint16_t { using IterationSteps = o2::utils::EnumFlags; struct TrackingParameters { + LayerMask getSeedingLayerMask() const noexcept + { + const auto layerSpan = LayerMask::span(0, NLayers - 1); + return SeedingLayers.empty() ? layerSpan : (SeedingLayers & layerSpan); + } + + int getNSeedingLayers() const noexcept + { + return getSeedingLayerMask().count(); + } + int CellMinimumLevel() const noexcept { const int minClusters = MinTrackLength - (MaxHoles > 0 ? MaxHoles : 0); const int effectiveMinClusters = minClusters > constants::ClustersPerCell ? minClusters : constants::ClustersPerCell; return effectiveMinClusters - constants::ClustersPerCell + 1; } - int NeighboursPerRoad() const noexcept { return NLayers - 3; } - int CellsPerRoad() const noexcept { return NLayers - 2; } - int TrackletsPerRoad() const noexcept { return NLayers - 1; } + int NeighboursPerRoad() const noexcept { return getNSeedingLayers() - 3; } + int CellsPerRoad() const noexcept { return getNSeedingLayers() - 2; } + int TrackletsPerRoad() const noexcept { return getNSeedingLayers() - 1; } std::string asString() const; IterationSteps PassFlags{IterationStep::FirstPass, IterationStep::RebuildClusterLUT}; @@ -76,6 +87,7 @@ struct TrackingParameters { int MinTrackLength = 7; int MaxHoles = 0; LayerMask HoleLayerMask = 0; + LayerMask SeedingLayers = 0; float NSigmaCut = 5; float PVres = 1.e-2f; /// Trackleting cuts diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h index 28bf4051fcc5f..dff6941d8c116 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h @@ -21,6 +21,7 @@ #include "ITStracking/Cell.h" #include "ITStracking/Cluster.h" #include "ITStracking/Constants.h" +#include "ITStracking/LayerMask.h" #include "ITStracking/MathUtils.h" #include "ITStracking/TrackITSInternal.h" #include "DetectorsBase/Propagator.h" @@ -45,6 +46,37 @@ GPUhdi() bool isBetter(const o2::its::TrackITSInternal& a, const o2::it return isBetter(a.getNumberOfClusters(), a.getChi2(), b.getNumberOfClusters(), b.getChi2()); } +GPUhdi() int getEffectiveTrackLength(LayerMask hitLayerMask, LayerMask nonCountingLayerMask) +{ + if (hitLayerMask.empty()) { + return 0; + } + return hitLayerMask.length() - (LayerMask::span(hitLayerMask.first(), hitLayerMask.last()) & nonCountingLayerMask).count(); +} + +GPUhdi() LayerMask getEffectiveHoleMask(LayerMask hitLayerMask, LayerMask nonCountingLayerMask) +{ + return hitLayerMask.holeMask() & ~nonCountingLayerMask; +} + +template +struct TrackSeedSelector { + float maxQ2Pt; + float maxChi2; + int maxHoles; + int minTrackLength; + LayerMask holeLayerMask; + LayerMask nonCountingLayerMask; + + GPUhd() bool operator()(const TrackSeed& seed) const + { + const auto hitLayerMask = seed.getHitLayerMask(); + return !(seed.getQ2Pt() > maxQ2Pt || seed.getChi2() > maxChi2) && + getEffectiveTrackLength(hitLayerMask, nonCountingLayerMask) >= minTrackLength && + getEffectiveHoleMask(hitLayerMask, nonCountingLayerMask).isAllowedHoleMask(maxHoles, holeLayerMask); + } +}; + // Find the populated interior layer closest to the radial midpoint. // If no layer can be found, return constants::UnusedIndex. // Should minimize the sagitta bias. diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h index a9b240d2ef037..282a366f102a7 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h @@ -104,15 +104,19 @@ class TrackingTopology #endif }; - void init(int maxLayers, int maxHoles, Mask holeLayerMask) + void init(int maxLayers, int maxHoles, Mask holeLayerMask, Mask seedingLayerMask = 0) { clear(); mMaxLayers = o2::gpu::CAMath::Max(0, o2::gpu::CAMath::Min(maxLayers, NLayers)); mMaxHoles = o2::gpu::CAMath::Max(maxHoles, 0); mHoleLayerMask = holeLayerMask; + mSeedingLayerMask = seedingLayerMask.empty() ? Mask::span(0, mMaxLayers - 1) : (seedingLayerMask & Mask::span(0, mMaxLayers - 1)); for (int fromLayer = 0; fromLayer < mMaxLayers; ++fromLayer) { + if (!mSeedingLayerMask.has(fromLayer)) { + continue; + } for (int toLayer = fromLayer + 1; toLayer < mMaxLayers; ++toLayer) { - if (Mask::skipped(fromLayer, toLayer).isAllowedHoleMask(mMaxHoles, mHoleLayerMask)) { + if (mSeedingLayerMask.has(toLayer) && isAllowedSeedingLink(fromLayer, toLayer)) { mLinks[mNLinks++] = LayerLink{static_cast(fromLayer), static_cast(toLayer)}; } } @@ -126,7 +130,7 @@ class TrackingTopology continue; } const Mask hitMask{first.fromLayer, first.toLayer, second.toLayer}; - if (hitMask.isAllowed(mMaxHoles, mHoleLayerMask)) { + if ((hitMask.holeMask() & mSeedingLayerMask).isAllowedHoleMask(mMaxHoles, mHoleLayerMask)) { mCells[mNCells++] = CellTopology{firstId, secondId, hitMask}; } } @@ -202,9 +206,15 @@ class TrackingTopology mNCellsByFirstLink = offset; } + bool isAllowedSeedingLink(int fromLayer, int toLayer) const noexcept + { + return (Mask::skipped(fromLayer, toLayer) & mSeedingLayerMask).isAllowedHoleMask(mMaxHoles, mHoleLayerMask); + } + int mMaxLayers{0}; int mMaxHoles{0}; Mask mHoleLayerMask{0}; + Mask mSeedingLayerMask{0}; Id mNLinks{0}; Id mNCells{0}; Id mNCellsByFirstLink{0}; diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx index 3bf4580d32157..27bef96621653 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx @@ -59,6 +59,9 @@ std::string TrackingParameters::asString() const if (MaxHoles) { str += std::format(" MaxHoles:{} HoleMask:{}", MaxHoles, HoleLayerMask.asString()); } + if (!SeedingLayers.empty()) { + str += std::format(" SeedingLayers:{}", SeedingLayers.asString()); + } if (PassFlags[IterationStep::TrackFollowerTop] || PassFlags[IterationStep::TrackFollowerBot]) { const bool top = PassFlags[IterationStep::TrackFollowerTop], bot = PassFlags[IterationStep::TrackFollowerBot]; str += std::format(" TrackFollower:{} NSigmaZ/Phi:{:.2f}/{:.2f}", diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index d418911202991..4477c8c33f699 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -249,7 +249,7 @@ void TimeFrame::initVertexingTopology(const TrackingParameters& trkPara template void TimeFrame::initDefaultTrackingTopology(const TrackingParameters& trkParam, const int maxLayers) { - mDefaultTrackingTopology.init(maxLayers, trkParam.MaxHoles, trkParam.HoleLayerMask); + mDefaultTrackingTopology.init(maxLayers, trkParam.MaxHoles, trkParam.HoleLayerMask, trkParam.getSeedingLayerMask()); } template @@ -258,7 +258,7 @@ void TimeFrame::initTrackerTopologies(gsl::span::findRoads(const int iteration) unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } const auto topology = mTimeFrame->getTrackingTopologyView(); + const auto nonCountingLayers = ~mTrkParams[iteration].getSeedingLayerMask(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { - auto seedFilter = [&](const auto& seed) { - return seed.getHitLayerMask().isAllowed(mTrkParams[iteration].MaxHoles, mTrkParams[iteration].HoleLayerMask) && - seed.getHitLayerMask().length() >= mTrkParams[iteration].MinTrackLength && - seed.getQ2Pt() <= 1.e3 && seed.getChi2() <= mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5); - }; + const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, mTrkParams[iteration].MinTrackLength, mTrkParams[iteration].HoleLayerMask, nonCountingLayers}; bounded_vector trackSeeds(mMemoryPool.get()); for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) { From fc218b979dc83bbc7206e80b5c8f64f8cd6d8e32 Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Tue, 9 Jun 2026 20:03:44 +0200 Subject: [PATCH 2/5] ITS: distinguish inactive layers from track holes Add an InactiveLayerMask to TrackingParameters to represent detector layers that are absent for a given input or configuration. Inactive layers remain skippable, but they do not inflate the effective seed length or consume the ordinary hole budget. Guard the multiple-scattering helper against zero material so inactive layers can be configured with no material budget. --- Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx | 2 +- .../ITSMFT/ITS/tracking/include/ITStracking/Configuration.h | 1 + Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h | 3 +++ Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx | 3 +++ Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 2 +- 5 files changed, 9 insertions(+), 2 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 4d43b765ef325..d9eff4dd20922 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -308,7 +308,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; const bool extendTracks = extendTop || extendBot; - const auto nonCountingLayers = ~this->mTrkParams[iteration].getSeedingLayerMask(); + const auto nonCountingLayers = this->mTrkParams[iteration].InactiveLayerMask | ~this->mTrkParams[iteration].getSeedingLayerMask(); for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { bounded_vector> trackSeeds(this->getMemoryPool().get()); for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) { diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 61ca263230bb9..83cb09b8ba439 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -87,6 +87,7 @@ struct TrackingParameters { int MinTrackLength = 7; int MaxHoles = 0; LayerMask HoleLayerMask = 0; + LayerMask InactiveLayerMask = 0; LayerMask SeedingLayers = 0; float NSigmaCut = 5; float PVres = 1.e-2f; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h index bff0b742e4547..5599755c52e64 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h @@ -122,6 +122,9 @@ GPUhdi() constexpr float SqDiff(float x, float y) GPUhdi() float MSangle(float mass, float p, float xX0) { + if (xX0 <= 0.f) { + return 0.f; + } float beta = p / o2::gpu::CAMath::Hypot(mass, p); return 0.0136f * o2::gpu::CAMath::Sqrt(xX0) * (1.f + 0.038f * o2::gpu::CAMath::Log(xX0)) / (beta * p); } diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx index 27bef96621653..eb4d90ee9d15f 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx @@ -59,6 +59,9 @@ std::string TrackingParameters::asString() const if (MaxHoles) { str += std::format(" MaxHoles:{} HoleMask:{}", MaxHoles, HoleLayerMask.asString()); } + if (!InactiveLayerMask.empty()) { + str += std::format(" InactiveMask:{}", InactiveLayerMask.asString()); + } if (!SeedingLayers.empty()) { str += std::format(" SeedingLayers:{}", SeedingLayers.asString()); } diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index a02f270c36a6c..1c29c54fa2c83 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -759,7 +759,7 @@ void TrackerTraits::findRoads(const int iteration) unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } const auto topology = mTimeFrame->getTrackingTopologyView(); - const auto nonCountingLayers = ~mTrkParams[iteration].getSeedingLayerMask(); + const auto nonCountingLayers = mTrkParams[iteration].InactiveLayerMask | ~mTrkParams[iteration].getSeedingLayerMask(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, mTrkParams[iteration].MinTrackLength, mTrkParams[iteration].HoleLayerMask, nonCountingLayers}; From f87bce2ff2beee751d1693f3db54055926c90285 Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Wed, 10 Jun 2026 00:14:22 +0200 Subject: [PATCH 3/5] ITS: make sure that seeds are not cut by min track lenght --- .../GPU/ITStrackingGPU/TrackingKernels.h | 2 +- .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 7 ++--- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 8 +++--- .../include/ITStracking/Configuration.h | 26 +++++++++++++++---- .../include/ITStracking/TrackingTopology.h | 5 ++++ .../ITSMFT/ITS/tracking/src/TimeFrame.cxx | 15 ++++++++--- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 14 ++++++++-- 7 files changed, 59 insertions(+), 18 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index e0d53766f08d3..9cefa59a14705 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -193,7 +193,7 @@ void processNeighboursHandler(const int startLevel, const float MaxChi2ClusterAttachment, const float maxChi2NDF, const int maxHoles, - const int minTrackLength, + const int minSeedingClusters, const LayerMask holeLayerMask, const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index d9eff4dd20922..98acca5105d80 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -308,7 +308,8 @@ void TrackerTraitsGPU::findRoads(const int iteration) const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; const bool extendTracks = extendTop || extendBot; - const auto nonCountingLayers = this->mTrkParams[iteration].InactiveLayerMask | ~this->mTrkParams[iteration].getSeedingLayerMask(); + const auto notSeedingLayers = this->mTrkParams[iteration].getNotSeedingLayerMask(); + const auto minSeedingClusters = this->mTrkParams[iteration].getMinSeedingClusters(); for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { bounded_vector> trackSeeds(this->getMemoryPool().get()); for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) { @@ -332,9 +333,9 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].MaxChi2NDF, this->mTrkParams[iteration].MaxHoles, - this->mTrkParams[iteration].MinTrackLength, + minSeedingClusters, this->mTrkParams[iteration].HoleLayerMask, - nonCountingLayers, + notSeedingLayers, this->mTrkParams[iteration].LayerxX0, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[iteration].CorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 1f974bfa3cda2..c1f7681b68880 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -958,7 +958,7 @@ void processNeighboursHandler(const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int maxHoles, - const int minTrackLength, + const int minSeedingClusters, const LayerMask holeLayerMask, const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, @@ -1094,7 +1094,7 @@ void processNeighboursHandler(const int startLevel, } GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream)); thrust::device_vector, gpu::TypedAllocator>> outSeeds(updatedCellSeed.size(), allocTrackSeed); - auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{1.e3f, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minTrackLength, holeLayerMask, nonCountingLayerMask}); + auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{1.e3f, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minSeedingClusters, holeLayerMask, nonCountingLayerMask}); auto s{end - outSeeds.begin()}; seedsHost.reserve(seedsHost.size() + s); thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost)); @@ -1364,7 +1364,7 @@ template void processNeighboursHandler<7>(const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int maxHoles, - const int minTrackLength, + const int minSeedingClusters, const LayerMask holeLayerMask, const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, @@ -1565,7 +1565,7 @@ template void processNeighboursHandler<11>(const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int maxHoles, - const int minTrackLength, + const int minSeedingClusters, const LayerMask holeLayerMask, const LayerMask nonCountingLayerMask, const std::vector& layerxX0Host, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 83cb09b8ba439..3ab070499f1dd 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -46,10 +46,20 @@ enum class IterationStep : uint16_t { using IterationSteps = o2::utils::EnumFlags; struct TrackingParameters { + LayerMask getActiveLayerMask() const noexcept + { + return LayerMask::span(0, NLayers - 1) & ~InactiveLayerMask; + } + LayerMask getSeedingLayerMask() const noexcept { - const auto layerSpan = LayerMask::span(0, NLayers - 1); - return SeedingLayers.empty() ? layerSpan : (SeedingLayers & layerSpan); + const auto activeLayers = getActiveLayerMask(); + return SeedingLayers.empty() ? activeLayers : (SeedingLayers & activeLayers); + } + + LayerMask getNotSeedingLayerMask() const noexcept + { + return ~getSeedingLayerMask(); } int getNSeedingLayers() const noexcept @@ -57,11 +67,17 @@ struct TrackingParameters { return getSeedingLayerMask().count(); } - int CellMinimumLevel() const noexcept + int getMinSeedingClusters() const noexcept { const int minClusters = MinTrackLength - (MaxHoles > 0 ? MaxHoles : 0); - const int effectiveMinClusters = minClusters > constants::ClustersPerCell ? minClusters : constants::ClustersPerCell; - return effectiveMinClusters - constants::ClustersPerCell + 1; + const int minClustersWithCells = minClusters > constants::ClustersPerCell ? minClusters : constants::ClustersPerCell; + const int nSeedingLayers = getNSeedingLayers(); + return minClustersWithCells < nSeedingLayers ? minClustersWithCells : nSeedingLayers; + } + + int CellMinimumLevel() const noexcept + { + return getMinSeedingClusters() - constants::ClustersPerCell + 1; } int NeighboursPerRoad() const noexcept { return getNSeedingLayers() - 3; } int CellsPerRoad() const noexcept { return getNSeedingLayers() - 2; } diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h index 282a366f102a7..5ee44c40d5776 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h @@ -111,6 +111,11 @@ class TrackingTopology mMaxHoles = o2::gpu::CAMath::Max(maxHoles, 0); mHoleLayerMask = holeLayerMask; mSeedingLayerMask = seedingLayerMask.empty() ? Mask::span(0, mMaxLayers - 1) : (seedingLayerMask & Mask::span(0, mMaxLayers - 1)); +#ifndef GPUCA_GPUCODE + if (mSeedingLayerMask.count() < constants::ClustersPerCell) { + LOGP(fatal, "Tracking topology has {} seeding layers, but at least {} are required to build CA cells", mSeedingLayerMask.count(), constants::ClustersPerCell); + } +#endif for (int fromLayer = 0; fromLayer < mMaxLayers; ++fromLayer) { if (!mSeedingLayerMask.has(fromLayer)) { continue; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index 4477c8c33f699..1f0d63abdf598 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -249,7 +249,10 @@ void TimeFrame::initVertexingTopology(const TrackingParameters& trkPara template void TimeFrame::initDefaultTrackingTopology(const TrackingParameters& trkParam, const int maxLayers) { - mDefaultTrackingTopology.init(maxLayers, trkParam.MaxHoles, trkParam.HoleLayerMask, trkParam.getSeedingLayerMask()); + if (maxLayers < trkParam.NLayers) { + LOGP(fatal, "Default tracking topology limited to {} layers, but the tracking parameters expect {}", maxLayers, trkParam.NLayers); + } + mDefaultTrackingTopology.init(trkParam.NLayers, trkParam.MaxHoles, trkParam.HoleLayerMask, trkParam.getSeedingLayerMask()); } template @@ -257,8 +260,14 @@ void TimeFrame::initTrackerTopologies(gsl::span nActiveLayers) { + LOGP(fatal, "Iteration {}: MinTrackLength {} cannot be satisfied with {} active layers", iteration, trkParams[iteration].MinTrackLength, nActiveLayers); + } + mTrackerTopologies[iteration].init(trkParams[iteration].NLayers, trkParams[iteration].MaxHoles, trkParams[iteration].HoleLayerMask, trkParams[iteration].getSeedingLayerMask()); } } diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 1c29c54fa2c83..032304acf5d28 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -759,10 +759,11 @@ void TrackerTraits::findRoads(const int iteration) unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } const auto topology = mTimeFrame->getTrackingTopologyView(); - const auto nonCountingLayers = mTrkParams[iteration].InactiveLayerMask | ~mTrkParams[iteration].getSeedingLayerMask(); + const auto notSeedingLayers = mTrkParams[iteration].getNotSeedingLayerMask(); + const auto minSeedingClusters = mTrkParams[iteration].getMinSeedingClusters(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { - const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, mTrkParams[iteration].MinTrackLength, mTrkParams[iteration].HoleLayerMask, nonCountingLayers}; + const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, minSeedingClusters, mTrkParams[iteration].HoleLayerMask, notSeedingLayers}; bounded_vector trackSeeds(mMemoryPool.get()); for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) { @@ -852,14 +853,18 @@ void TrackerTraits::acceptTracks(int iteration, auto& trks = mTimeFrame->getTracks(); trks.reserve(trks.size() + tracks.size()); const float smallestROFHalf = mTimeFrame->getROFOverlapTableView().getClockLayer().mROFLength * 0.5f; + const int minTrackLength = mTrkParams[iteration].MinTrackLength; + const LayerMask inactiveLayerMask = mTrkParams[iteration].InactiveLayerMask; for (auto& track : tracks) { int nShared = 0; bool isFirstShared{false}; int firstLayer{-1}, firstCluster{-1}; + LayerMask hitLayerMask{0}; for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { if (track.getClusterIndex(iLayer) == constants::UnusedIndex) { continue; } + hitLayerMask.set(iLayer); bool isShared = mTimeFrame->isClusterUsed(iLayer, track.getClusterIndex(iLayer)); nShared += int(isShared); if (firstLayer < 0) { @@ -869,6 +874,11 @@ void TrackerTraits::acceptTracks(int iteration, } } + /// seeds are selected with a length cut relaxed to the seeding layers: enforce the full minimum length before accepting the final track + if (track::getEffectiveTrackLength(hitLayerMask, inactiveLayerMask) < minTrackLength) { + continue; + } + /// do not account for the first cluster in the shared clusters number if it is allowed if (nShared - int(isFirstShared && mTrkParams[iteration].AllowSharingFirstCluster) > mTrkParams[iteration].SharedMaxClusters) { continue; From b74614f4a9ffdb29ff0d37586d4d88b3c159a72e Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Fri, 12 Jun 2026 16:21:08 +0200 Subject: [PATCH 4/5] ITS: implement Felix' comments --- .../GPU/ITStrackingGPU/TrackingKernels.h | 2 +- .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 4 +- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 8 ++-- .../include/ITStracking/Configuration.h | 4 +- .../include/ITStracking/TrackHelpers.h | 45 +++++++++---------- .../include/ITStracking/TrackingTopology.h | 5 ++- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 6 +-- 7 files changed, 38 insertions(+), 36 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 9cefa59a14705..f4a89d9d24d8f 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -195,7 +195,7 @@ void processNeighboursHandler(const int startLevel, const int maxHoles, const int minSeedingClusters, const LayerMask holeLayerMask, - const LayerMask nonCountingLayerMask, + const LayerMask nonSeedingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 98acca5105d80..51d817943d395 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -308,7 +308,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; const bool extendTracks = extendTop || extendBot; - const auto notSeedingLayers = this->mTrkParams[iteration].getNotSeedingLayerMask(); + const auto nonSeedingLayerMask = this->mTrkParams[iteration].getNonSeedingLayerMask(); const auto minSeedingClusters = this->mTrkParams[iteration].getMinSeedingClusters(); for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { bounded_vector> trackSeeds(this->getMemoryPool().get()); @@ -335,7 +335,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].MaxHoles, minSeedingClusters, this->mTrkParams[iteration].HoleLayerMask, - notSeedingLayers, + nonSeedingLayerMask, this->mTrkParams[iteration].LayerxX0, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[iteration].CorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index c1f7681b68880..963f91fdc6678 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -960,7 +960,7 @@ void processNeighboursHandler(const int startLevel, const int maxHoles, const int minSeedingClusters, const LayerMask holeLayerMask, - const LayerMask nonCountingLayerMask, + const LayerMask nonSeedingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, @@ -1094,7 +1094,7 @@ void processNeighboursHandler(const int startLevel, } GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream)); thrust::device_vector, gpu::TypedAllocator>> outSeeds(updatedCellSeed.size(), allocTrackSeed); - auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{1.e3f, maxChi2NDF * ((startLevel + 2) * 2 - 5), maxHoles, minSeedingClusters, holeLayerMask, nonCountingLayerMask}); + auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{1.e3f, maxChi2NDF, startLevel, maxHoles, minSeedingClusters, holeLayerMask, nonSeedingLayerMask}); auto s{end - outSeeds.begin()}; seedsHost.reserve(seedsHost.size() + s); thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost)); @@ -1366,7 +1366,7 @@ template void processNeighboursHandler<7>(const int startLevel, const int maxHoles, const int minSeedingClusters, const LayerMask holeLayerMask, - const LayerMask nonCountingLayerMask, + const LayerMask nonSeedingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, @@ -1567,7 +1567,7 @@ template void processNeighboursHandler<11>(const int startLevel, const int maxHoles, const int minSeedingClusters, const LayerMask holeLayerMask, - const LayerMask nonCountingLayerMask, + const LayerMask nonSeedingLayerMask, const std::vector& layerxX0Host, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 3ab070499f1dd..8e1c68fc31c6c 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -57,9 +57,9 @@ struct TrackingParameters { return SeedingLayers.empty() ? activeLayers : (SeedingLayers & activeLayers); } - LayerMask getNotSeedingLayerMask() const noexcept + LayerMask getNonSeedingLayerMask() const noexcept { - return ~getSeedingLayerMask(); + return ~(getSeedingLayerMask()); } int getNSeedingLayers() const noexcept diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h index dff6941d8c116..77636333d703f 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h @@ -35,30 +35,11 @@ GPUhdi() bool isBetter(const int nClustersA, const float chi2A, const int nClust return (nClustersA > nClustersB) || (nClustersA == nClustersB && chi2A < chi2B); } -GPUhdi() bool isBetter(const o2::its::TrackITS& a, const o2::its::TrackITS& b) +GPUhdi() bool isBetter(const auto& a, const auto& b) { return isBetter(a.getNumberOfClusters(), a.getChi2(), b.getNumberOfClusters(), b.getChi2()); } -template -GPUhdi() bool isBetter(const o2::its::TrackITSInternal& a, const o2::its::TrackITSInternal& b) -{ - return isBetter(a.getNumberOfClusters(), a.getChi2(), b.getNumberOfClusters(), b.getChi2()); -} - -GPUhdi() int getEffectiveTrackLength(LayerMask hitLayerMask, LayerMask nonCountingLayerMask) -{ - if (hitLayerMask.empty()) { - return 0; - } - return hitLayerMask.length() - (LayerMask::span(hitLayerMask.first(), hitLayerMask.last()) & nonCountingLayerMask).count(); -} - -GPUhdi() LayerMask getEffectiveHoleMask(LayerMask hitLayerMask, LayerMask nonCountingLayerMask) -{ - return hitLayerMask.holeMask() & ~nonCountingLayerMask; -} - template struct TrackSeedSelector { float maxQ2Pt; @@ -66,14 +47,32 @@ struct TrackSeedSelector { int maxHoles; int minTrackLength; LayerMask holeLayerMask; - LayerMask nonCountingLayerMask; + LayerMask nonSeedingLayerMask; + + GPUhd() TrackSeedSelector(float maxQ2Pt, float maxChi2NDF, int startLevel, int maxHoles, int minTrackLength, LayerMask holeLayerMask, LayerMask nonSeedingLayerMask) + : maxQ2Pt{maxQ2Pt}, maxChi2{maxChi2NDF * ((startLevel + 2) * 2 - 5)}, maxHoles{maxHoles}, minTrackLength{minTrackLength}, holeLayerMask{holeLayerMask}, nonSeedingLayerMask{nonSeedingLayerMask} + { + } + + static GPUhdi() int getEffectiveTrackLength(LayerMask hitLayerMask, LayerMask excludedLayerMask) + { + if (hitLayerMask.empty()) { + return 0; + } + return hitLayerMask.length() - (LayerMask::span(hitLayerMask.first(), hitLayerMask.last()) & excludedLayerMask).count(); + } + + static GPUhdi() LayerMask getEffectiveHoleMask(LayerMask hitLayerMask, LayerMask excludedLayerMask) + { + return hitLayerMask.holeMask() & ~excludedLayerMask; + } GPUhd() bool operator()(const TrackSeed& seed) const { const auto hitLayerMask = seed.getHitLayerMask(); return !(seed.getQ2Pt() > maxQ2Pt || seed.getChi2() > maxChi2) && - getEffectiveTrackLength(hitLayerMask, nonCountingLayerMask) >= minTrackLength && - getEffectiveHoleMask(hitLayerMask, nonCountingLayerMask).isAllowedHoleMask(maxHoles, holeLayerMask); + getEffectiveTrackLength(hitLayerMask, nonSeedingLayerMask) >= minTrackLength && + getEffectiveHoleMask(hitLayerMask, nonSeedingLayerMask).isAllowedHoleMask(maxHoles, holeLayerMask); } }; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h index 5ee44c40d5776..f61c3e6c77c80 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h @@ -70,6 +70,7 @@ class TrackingTopology const CellTopology* cells{nullptr}; const Range* cellsByFirstLinkIndex{nullptr}; const Id* cellsByFirstLink{nullptr}; + Mask seedingLayerMask{0}; Id nLinks{0}; Id nCells{0}; Id nCellsByFirstLink{0}; @@ -81,7 +82,7 @@ class TrackingTopology #ifndef GPUCA_GPUCODE std::string asString() const { - std::string out = fmt::format("TrackingTopology: links={} cells={}", nLinks, nCells); + std::string out = fmt::format("TrackingTopology: links={} cells={} seedingLayers={}", nLinks, nCells, seedingLayerMask.asString()); out += "\n links:"; for (Id linkId = 0; linkId < nLinks; ++linkId) { const auto& t = links[linkId]; @@ -150,6 +151,7 @@ class TrackingTopology mCells.data(), mCellsByFirstLinkIndex.data(), mCellsByFirstLink.data(), + mSeedingLayerMask, mNLinks, mNCells, mNCellsByFirstLink}; @@ -164,6 +166,7 @@ class TrackingTopology deviceCells, deviceCellsByFirstLinkIndex, deviceCellsByFirstLink, + mSeedingLayerMask, mNLinks, mNCells, mNCellsByFirstLink}; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 032304acf5d28..d32965b9a68a0 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -759,11 +759,11 @@ void TrackerTraits::findRoads(const int iteration) unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } const auto topology = mTimeFrame->getTrackingTopologyView(); - const auto notSeedingLayers = mTrkParams[iteration].getNotSeedingLayerMask(); + const auto nonSeedingLayerMask = mTrkParams[iteration].getNonSeedingLayerMask(); const auto minSeedingClusters = mTrkParams[iteration].getMinSeedingClusters(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { - const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF * ((startLevel + 2) * 2 - 5), mTrkParams[iteration].MaxHoles, minSeedingClusters, mTrkParams[iteration].HoleLayerMask, notSeedingLayers}; + const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF, startLevel, mTrkParams[iteration].MaxHoles, minSeedingClusters, mTrkParams[iteration].HoleLayerMask, nonSeedingLayerMask}; bounded_vector trackSeeds(mMemoryPool.get()); for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) { @@ -875,7 +875,7 @@ void TrackerTraits::acceptTracks(int iteration, } /// seeds are selected with a length cut relaxed to the seeding layers: enforce the full minimum length before accepting the final track - if (track::getEffectiveTrackLength(hitLayerMask, inactiveLayerMask) < minTrackLength) { + if (track::TrackSeedSelector::getEffectiveTrackLength(hitLayerMask, inactiveLayerMask) < minTrackLength) { continue; } From 11b2171255241c39ee712729a0086a8ae6768cef Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Fri, 12 Jun 2026 21:29:01 +0200 Subject: [PATCH 5/5] ITS: second round of comments --- .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 6 ++--- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 2 +- .../tracking/include/ITStracking/Constants.h | 1 + .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 26 +++++++++---------- 4 files changed, 16 insertions(+), 19 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 51d817943d395..47d0ae0e34801 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -308,8 +308,6 @@ void TrackerTraitsGPU::findRoads(const int iteration) const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; const bool extendTracks = extendTop || extendBot; - const auto nonSeedingLayerMask = this->mTrkParams[iteration].getNonSeedingLayerMask(); - const auto minSeedingClusters = this->mTrkParams[iteration].getMinSeedingClusters(); for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { bounded_vector> trackSeeds(this->getMemoryPool().get()); for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) { @@ -333,9 +331,9 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].MaxChi2NDF, this->mTrkParams[iteration].MaxHoles, - minSeedingClusters, + this->mTrkParams[iteration].getMinSeedingClusters(), this->mTrkParams[iteration].HoleLayerMask, - nonSeedingLayerMask, + this->mTrkParams[iteration].getNonSeedingLayerMask(), this->mTrkParams[iteration].LayerxX0, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[iteration].CorrType, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 963f91fdc6678..6f4f2bca76722 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -1094,7 +1094,7 @@ void processNeighboursHandler(const int startLevel, } GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream)); thrust::device_vector, gpu::TypedAllocator>> outSeeds(updatedCellSeed.size(), allocTrackSeed); - auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{1.e3f, maxChi2NDF, startLevel, maxHoles, minSeedingClusters, holeLayerMask, nonSeedingLayerMask}); + auto end = thrust::copy_if(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), track::TrackSeedSelector{constants::MaxTrackSeedQ2Pt, maxChi2NDF, startLevel, maxHoles, minSeedingClusters, holeLayerMask, nonSeedingLayerMask}); auto s{end - outSeeds.begin()}; seedsHost.reserve(seedsHost.size() + s); thrust::copy(outSeeds.begin(), outSeeds.begin() + s, std::back_inserter(seedsHost)); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h index 0378de833d6eb..dfbaab4601e85 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h @@ -29,6 +29,7 @@ constexpr bool DoTimeBenchmarks = true; constexpr bool SaveTimeBenchmarks = false; constexpr float Tolerance = 1e-12; // numerical tolerance constexpr int ClustersPerCell = 3; // number of clusters for a cell +constexpr float MaxTrackSeedQ2Pt = 1.e3f; // maximum q/pt for track seeds constexpr int UnusedIndex = -1; // global unused flag constexpr float UnsetValue = -999.f; // global unset value constexpr float Radl = 9.36f; // Radiation length of Si [cm] diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index d32965b9a68a0..181417fbe16a0 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -686,12 +686,21 @@ bool TrackerTraits::finaliseTrackSeed(const TrackSeedN& seed, trkParams.ReseedIfShorter)) { return false; } + const auto passesFinalLengthCut = [&trkParams](const TrackITSExt& candidate) { + LayerMask hitLayerMask{0}; + for (int iLayer{0}; iLayer < trkParams.NLayers; ++iLayer) { + if (candidate.getClusterIndex(iLayer) != constants::UnusedIndex) { + hitLayerMask.set(iLayer); + } + } + return track::TrackSeedSelector::getEffectiveTrackLength(hitLayerMask, trkParams.InactiveLayerMask) >= trkParams.MinTrackLength; + }; const bool extendTop = trkParams.PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = trkParams.PassFlags[IterationStep::TrackFollowerBot]; if (!extendTop && !extendBot) { track = makeTrackITSExt(internalTrack); - return true; + return passesFinalLengthCut(track); } const int maxHypotheses = std::max(1, trkParams.TrackFollowerMaxHypotheses); @@ -743,7 +752,7 @@ bool TrackerTraits::finaliseTrackSeed(const TrackSeedN& seed, if (bestDiff) { track.setExtendedLayerPattern(bestDiff); } - return true; + return passesFinalLengthCut(track); } template @@ -759,11 +768,9 @@ void TrackerTraits::findRoads(const int iteration) unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } const auto topology = mTimeFrame->getTrackingTopologyView(); - const auto nonSeedingLayerMask = mTrkParams[iteration].getNonSeedingLayerMask(); - const auto minSeedingClusters = mTrkParams[iteration].getMinSeedingClusters(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { - const track::TrackSeedSelector seedFilter{1.e3f, mTrkParams[iteration].MaxChi2NDF, startLevel, mTrkParams[iteration].MaxHoles, minSeedingClusters, mTrkParams[iteration].HoleLayerMask, nonSeedingLayerMask}; + const track::TrackSeedSelector seedFilter{constants::MaxTrackSeedQ2Pt, mTrkParams[iteration].MaxChi2NDF, startLevel, mTrkParams[iteration].MaxHoles, mTrkParams[iteration].getMinSeedingClusters(), mTrkParams[iteration].HoleLayerMask, mTrkParams[iteration].getNonSeedingLayerMask()}; bounded_vector trackSeeds(mMemoryPool.get()); for (int startCellTopologyId{0}; startCellTopologyId < topology.nCells; ++startCellTopologyId) { @@ -853,18 +860,14 @@ void TrackerTraits::acceptTracks(int iteration, auto& trks = mTimeFrame->getTracks(); trks.reserve(trks.size() + tracks.size()); const float smallestROFHalf = mTimeFrame->getROFOverlapTableView().getClockLayer().mROFLength * 0.5f; - const int minTrackLength = mTrkParams[iteration].MinTrackLength; - const LayerMask inactiveLayerMask = mTrkParams[iteration].InactiveLayerMask; for (auto& track : tracks) { int nShared = 0; bool isFirstShared{false}; int firstLayer{-1}, firstCluster{-1}; - LayerMask hitLayerMask{0}; for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { if (track.getClusterIndex(iLayer) == constants::UnusedIndex) { continue; } - hitLayerMask.set(iLayer); bool isShared = mTimeFrame->isClusterUsed(iLayer, track.getClusterIndex(iLayer)); nShared += int(isShared); if (firstLayer < 0) { @@ -874,11 +877,6 @@ void TrackerTraits::acceptTracks(int iteration, } } - /// seeds are selected with a length cut relaxed to the seeding layers: enforce the full minimum length before accepting the final track - if (track::TrackSeedSelector::getEffectiveTrackLength(hitLayerMask, inactiveLayerMask) < minTrackLength) { - continue; - } - /// do not account for the first cluster in the shared clusters number if it is allowed if (nShared - int(isFirstShared && mTrkParams[iteration].AllowSharingFirstCluster) > mTrkParams[iteration].SharedMaxClusters) { continue;