From 5b07ba327dd9f7f3ca24a9c6a04bcf6c4595bdc9 Mon Sep 17 00:00:00 2001 From: Felix Schlepper Date: Sat, 25 Apr 2026 15:13:21 +0200 Subject: [PATCH 1/6] 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 --- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 2 +- .../include/ITStracking/IndexTableUtils.h | 2 +- .../include/ITStracking/VertexerTraits.h | 34 ------------------- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 2 +- .../ITS/tracking/src/VertexerTraits.cxx | 4 +-- 5 files changed, 5 insertions(+), 39 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 4b12583d99c00..7cf0ef4c0b0d7 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -408,7 +408,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel( const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + constants::Tolerance)}; /// protecting from overflows adding the detector resolution 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))}; const int4 selectedBinsRect{o2::its::getBinsRect(currentCluster, layerIndex + 1, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut, *utils)}; - if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) { + if (selectedBinsRect.x < 0) { continue; } int phiBinsNum{selectedBinsRect.w - selectedBinsRect.y + 1}; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h index e2487208e9453..4e8d5bcfea42a 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h @@ -124,7 +124,7 @@ GPUhdi() int4 getBinsRect(const Cluster& currentCluster, const int layerIndex, if (zRangeMax < -utils.getLayerZ(layerIndex) || zRangeMin > utils.getLayerZ(layerIndex) || zRangeMin > zRangeMax) { - return int4{0, 0, 0, 0}; + return int4{-1, -1, -1, -1}; } return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex, zRangeMin)), diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h index 1adb09551e326..3230737a0f87c 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h @@ -53,12 +53,6 @@ class VertexerTraits VertexerTraits() = default; virtual ~VertexerTraits() = default; - GPUhdi() static consteval int4 getEmptyBinsRect() - { - return int4{0, 0, 0, 0}; - } - GPUhd() const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi); - GPUhd() static const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi, const IndexTableUtilsN&); GPUhd() static const int2 getPhiBins(float phi, float deltaPhi, const IndexTableUtilsN&); GPUhd() const int2 getPhiBins(float phi, float deltaPhi) { return getPhiBins(phi, deltaPhi, mIndexTableUtils); } @@ -134,34 +128,6 @@ GPUhdi() const int2 VertexerTraits::getPhiBins(float phi, float dPhi, c utils.getPhiBinIndex(math_utils::getNormalizedPhi(phi + dPhi))}; } -template -GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, const int layerIndex, - const float directionZIntersection, float maxdeltaz, float maxdeltaphi, - const IndexTableUtilsN& utils) -{ - const float zRangeMin = directionZIntersection - 2 * maxdeltaz; - const float phiRangeMin = currentCluster.phi - maxdeltaphi; - const float zRangeMax = directionZIntersection + 2 * maxdeltaz; - const float phiRangeMax = currentCluster.phi + maxdeltaphi; - - if (zRangeMax < -utils.getLayerZ(layerIndex + 1) || - zRangeMin > utils.getLayerZ(layerIndex + 1) || zRangeMin > zRangeMax) { - return getEmptyBinsRect(); - } - - return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex + 1, zRangeMin)), - utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMin)), - o2::gpu::GPUCommonMath::Min(utils.getNzBins() - 1, utils.getZBinIndex(layerIndex + 1, zRangeMax)), - utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))}; -} - -template -GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, const int layerIndex, - const float directionZIntersection, float maxdeltaz, float maxdeltaphi) -{ - return VertexerTraits::getBinsRect(currentCluster, layerIndex, directionZIntersection, maxdeltaz, maxdeltaphi, mIndexTableUtils); -} - } // namespace its } // namespace o2 #endif diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 69ccd8228ad88..b66e2f7240bfb 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -114,7 +114,7 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer const auto bins = o2::its::getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax, sigmaZ * mTrkParams[iteration].NSigmaCut, mTimeFrame->getPhiCut(iLayer), mTimeFrame->getIndexTableUtils()); - if (bins.x == 0 && bins.y == 0 && bins.z == 0 && bins.w == 0) { + if (bins.x < 0) { continue; } int phiBinsNum = bins.w - bins.y + 1; diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index d0baa65c49147..5a2781839f035 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -58,8 +58,8 @@ void trackleterKernelHost( for (int iCurrentLayerClusterIndex = 0; iCurrentLayerClusterIndex < clustersCurrentLayer.size(); ++iCurrentLayerClusterIndex) { int storedTracklets{0}; const Cluster& currentCluster{clustersCurrentLayer[iCurrentLayerClusterIndex]}; - const int4 selectedBinsRect{VertexerTraits::getBinsRect(currentCluster, (int)Mode, 0.f, 50.f, phiCut / 2, utils)}; - if (selectedBinsRect.x != 0 || selectedBinsRect.y != 0 || selectedBinsRect.z != 0 || selectedBinsRect.w != 0) { + const int4 selectedBinsRect{o2::its::getBinsRect(currentCluster, (int)Mode + 1, 0.f, 0.f, 100.f, phiCut / 2, utils)}; + if (selectedBinsRect.x >= 0) { int phiBinsNum{selectedBinsRect.w - selectedBinsRect.y + 1}; if (phiBinsNum < 0) { phiBinsNum += PhiBins; From a41296956f2a378a3e4d71ee3acd5cbfc7fd94da Mon Sep 17 00:00:00 2001 From: Felix Schlepper Date: Sat, 25 Apr 2026 16:11:09 +0200 Subject: [PATCH 2/6] ITS: speedup line selection Signed-off-by: Felix Schlepper --- .../ITS/tracking/include/ITStracking/Cell.h | 7 ++-- .../tracking/include/ITStracking/Constants.h | 16 ++++--- .../tracking/include/ITStracking/MathUtils.h | 6 +++ .../tracking/include/ITStracking/TimeFrame.h | 6 +-- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 3 +- .../ITS/tracking/src/VertexerTraits.cxx | 42 +++++++++++-------- 6 files changed, 43 insertions(+), 37 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h index dce66bdf99415..0cd81082eda32 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h @@ -64,10 +64,9 @@ class SeedBase : public o2::track::TrackParCovF }; /// CellSeed: connections of three clusters -class CellSeed final : public SeedBase<3> +class CellSeed final : public SeedBase { - static constexpr int NStoredClusters = 3; - using Base = SeedBase; + using Base = SeedBase; public: GPUhdDefault() CellSeed() = default; @@ -98,7 +97,7 @@ class CellSeed final : public SeedBase<3> GPUhd() int getCluster(int layer) const { const int rel = layer - getInnerLayer(); - return (rel >= 0 && rel < NStoredClusters) ? this->clustersRaw()[rel] : constants::UnusedIndex; + return (rel >= 0 && rel < constants::ClustersPerCell) ? this->clustersRaw()[rel] : constants::UnusedIndex; } }; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h index d48e8fb7c5856..85f4b3a222af2 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h @@ -19,9 +19,6 @@ #include #include -#include "GPUCommonDef.h" -#include "GPUCommonDefAPI.h" - namespace o2::its::constants { @@ -30,12 +27,13 @@ constexpr float MB = KB * KB; constexpr float GB = MB * KB; constexpr bool DoTimeBenchmarks = true; constexpr bool SaveTimeBenchmarks = false; -constexpr float Tolerance = 1e-12; // numerical tolerance -constexpr int ClustersPerCell = 3; -constexpr int UnusedIndex = -1; -constexpr float Radl = 9.36f; // Radiation length of Si [cm] -constexpr float Rho = 2.33f; // Density of Si [g/cm^3] -constexpr int MaxIter = 4; // Max. supported iterations +constexpr float Tolerance = 1e-12; // numerical tolerance +constexpr int ClustersPerCell = 3; // number of clusters for a cell +constexpr int UnusedIndex = -1; // global unused flag +constexpr float Radl = 9.36f; // Radiation length of Si [cm] +constexpr float Rho = 2.33f; // Density of Si [g/cm^3] +constexpr int MaxIter = 4; // Max. supported iterations +constexpr int MaxSelectedTrackletsPerCluster = 100; // vertexer: max lines per cluster namespace helpers { diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h index ab3c7d5d29873..d276e27638dbd 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/MathUtils.h @@ -89,6 +89,12 @@ GPUhdi() float smallestAngleDifference(float a, float b) return o2::gpu::CAMath::Remainderf(b - a, o2::constants::math::TwoPI); } +GPUhdi() bool isPhiDifferenceBelow(const float phiA, const float phiB, const float phiCut) +{ + const float deltaPhi = o2::gpu::CAMath::Abs(phiA - phiB); + return deltaPhi < phiCut || deltaPhi > o2::constants::math::TwoPI - phiCut; +} + GPUhdi() constexpr float Sq(float v) { return v * v; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index f2506694755c5..300abb2a3b10d 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -26,9 +26,7 @@ #include "ITStracking/Cell.h" #include "ITStracking/Cluster.h" #include "ITStracking/Configuration.h" -#include "ITStracking/Constants.h" #include "ITStracking/ClusterLines.h" -#include "ITStracking/Definitions.h" #include "ITStracking/Tracklet.h" #include "ITStracking/IndexTableUtils.h" #include "ITStracking/ExternalAllocator.h" @@ -103,7 +101,7 @@ struct TimeFrame { void setBeamPosition(const float x, const float y, const float s2, const float base = 50.f, const float systematic = 0.f) { isBeamPositionOverridden = true; - resetBeamXY(x, y, s2 / o2::gpu::CAMath::Sqrt(base * base + systematic)); + resetBeamXY(x, y, s2 / o2::gpu::CAMath::Sqrt((base * base) + systematic)); } float getBeamX() const { return mBeamPos[0]; } @@ -249,7 +247,7 @@ struct TimeFrame { // Propagator const o2::base::PropagatorImpl* getDevicePropagator() const { return mPropagatorDevice; } - virtual void setDevicePropagator(const o2::base::PropagatorImpl*) {}; + virtual void setDevicePropagator(const o2::base::PropagatorImpl* /*unused*/) {}; template void addClusterToLayer(int layer, T&&... args); diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index b66e2f7240bfb..d5db4b983f754 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -150,11 +150,10 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer if (mTimeFrame->isClusterUsed(iLayer + 1, nextCluster.clusterId)) { continue; } - const float deltaPhi = o2::gpu::CAMath::Abs(o2::math_utils::toPMPi(currentCluster.phi - nextCluster.phi)); const float deltaZ = o2::gpu::CAMath::Abs((tanLambda * (nextCluster.radius - currentCluster.radius)) + currentCluster.zCoordinate - nextCluster.zCoordinate); if (deltaZ / sigmaZ < mTrkParams[iteration].NSigmaCut && - ((deltaPhi < mTimeFrame->getPhiCut(iLayer) || o2::gpu::GPUCommonMath::Abs(deltaPhi - o2::constants::math::TwoPI) < mTimeFrame->getPhiCut(iLayer)))) { + math_utils::isPhiDifferenceBelow(currentCluster.phi, nextCluster.phi, mTimeFrame->getPhiCut(iLayer))) { const float phi{o2::gpu::CAMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate)}; const float tanL = (currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius); if constexpr (decltype(Tag)::value == PassMode::OnePass::value) { diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index 5a2781839f035..6e0b406cddfd1 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -36,6 +36,7 @@ namespace o2::its { namespace { + template void trackleterKernelHost( const gsl::span& clustersNextLayer, // 0 2 @@ -48,9 +49,9 @@ void trackleterKernelHost( const IndexTableUtils& utils, const TimeEstBC& timErr, gsl::span rofFoundTrackletsOffsets, - const int globalOffsetNextLayer = 0, - const int globalOffsetCurrentLayer = 0, - const int maxTrackletsPerCluster = static_cast(2e3)) + const int globalOffsetNextLayer, + const int globalOffsetCurrentLayer, + const int maxTrackletsPerCluster) { const int PhiBins{utils.getNphiBins()}; const int ZBins{utils.getNzBins()}; @@ -65,17 +66,17 @@ void trackleterKernelHost( phiBinsNum += PhiBins; } // loop on phi bins next layer - for (int iPhiBin{selectedBinsRect.y}, iPhiCount{0}; iPhiCount < phiBinsNum; iPhiBin = ++iPhiBin == PhiBins ? 0 : iPhiBin, iPhiCount++) { + for (int iPhiBin{selectedBinsRect.y}, iPhiCount{0}; iPhiCount < phiBinsNum && storedTracklets < maxTrackletsPerCluster; iPhiBin = ++iPhiBin == PhiBins ? 0 : iPhiBin, iPhiCount++) { const int firstBinIndex{utils.getBinIndex(selectedBinsRect.x, iPhiBin)}; const int firstRowClusterIndex{indexTableNext[firstBinIndex]}; const int maxRowClusterIndex{indexTableNext[firstBinIndex + ZBins]}; // loop on clusters next layer - for (int iNextLayerClusterIndex{firstRowClusterIndex}; iNextLayerClusterIndex < maxRowClusterIndex && iNextLayerClusterIndex < static_cast(clustersNextLayer.size()); ++iNextLayerClusterIndex) { + for (int iNextLayerClusterIndex{firstRowClusterIndex}; iNextLayerClusterIndex < maxRowClusterIndex && iNextLayerClusterIndex < static_cast(clustersNextLayer.size()) && storedTracklets < maxTrackletsPerCluster; ++iNextLayerClusterIndex) { if (usedClustersNextLayer[iNextLayerClusterIndex]) { continue; } const Cluster& nextCluster{clustersNextLayer[iNextLayerClusterIndex]}; - if (o2::gpu::GPUCommonMath::Abs(math_utils::smallestAngleDifference(currentCluster.phi, nextCluster.phi)) < phiCut) { + if (math_utils::isPhiDifferenceBelow(currentCluster.phi, nextCluster.phi, phiCut)) { if (storedTracklets < maxTrackletsPerCluster) { if constexpr (!EvalRun) { if constexpr (Mode == TrackletMode::Layer0Layer1) { @@ -105,35 +106,39 @@ void trackletSelectionKernelHost( gsl::span usedClusters2, // global layer 2 used clusters const gsl::span& tracklets01, const gsl::span& tracklets12, - bounded_vector& usedTracklets, + bounded_vector& usedTracklets, const gsl::span foundTracklets01, const gsl::span foundTracklets12, bounded_vector& lines, const gsl::span& trackletLabels, bounded_vector& linesLabels, const int nLayer1Clusters, - const float tanLambdaCut = 0.025f, - const float phiCut = 0.005f, - const int maxTracklets = 100) + const float tanLambdaCut, + const float phiCut, + const int maxTracklets) { int offset01{0}, offset12{0}; for (int iCurrentLayerClusterIndex{0}; iCurrentLayerClusterIndex < nLayer1Clusters; ++iCurrentLayerClusterIndex) { int validTracklets{0}; - for (int iTracklet12{offset12}; iTracklet12 < offset12 + foundTracklets12[iCurrentLayerClusterIndex]; ++iTracklet12) { - for (int iTracklet01{offset01}; iTracklet01 < offset01 + foundTracklets01[iCurrentLayerClusterIndex]; ++iTracklet01) { + const int endTracklet01 = offset01 + foundTracklets01[iCurrentLayerClusterIndex]; + const int endTracklet12 = offset12 + foundTracklets12[iCurrentLayerClusterIndex]; + for (int iTracklet12{offset12}; iTracklet12 < endTracklet12 && validTracklets != maxTracklets; ++iTracklet12) { + const auto& tracklet12{tracklets12[iTracklet12]}; + for (int iTracklet01{offset01}; iTracklet01 < endTracklet01 && validTracklets != maxTracklets; ++iTracklet01) { if (usedTracklets[iTracklet01]) { continue; } const auto& tracklet01{tracklets01[iTracklet01]}; - const auto& tracklet12{tracklets12[iTracklet12]}; if (!tracklet01.getTimeStamp().isCompatible(tracklet12.getTimeStamp())) { continue; } const float deltaTanLambda{o2::gpu::GPUCommonMath::Abs(tracklet01.tanLambda - tracklet12.tanLambda)}; - const float deltaPhi{o2::gpu::GPUCommonMath::Abs(math_utils::smallestAngleDifference(tracklet01.phi, tracklet12.phi))}; - if (deltaTanLambda < tanLambdaCut && deltaPhi < phiCut && validTracklets != maxTracklets) { + if (deltaTanLambda >= tanLambdaCut) { + continue; + } + if (math_utils::isPhiDifferenceBelow(tracklet01.phi, tracklet12.phi, phiCut) && validTracklets != maxTracklets) { usedClusters0[tracklet01.firstClusterIndex] = 1; usedClusters2[tracklet12.secondClusterIndex] = 1; usedTracklets[iTracklet01] = true; @@ -296,8 +301,8 @@ void VertexerTraits::computeTrackletMatching(const int iteration) if (mTimeFrame->getFoundTracklets(pivotRofId, 0).empty() || skipROF(iteration, pivotRofId)) { continue; } - mTimeFrame->getLines(pivotRofId).reserve(mTimeFrame->getNTrackletsCluster(pivotRofId, 0).size()); - bounded_vector usedTracklets(mTimeFrame->getFoundTracklets(pivotRofId, 0).size(), false, mMemoryPool.get()); + mTimeFrame->getLines(pivotRofId).reserve(std::min(mTimeFrame->getFoundTracklets(pivotRofId, 0).size(), mTimeFrame->getNTrackletsCluster(pivotRofId, 0).size() * constants::MaxSelectedTrackletsPerCluster)); + bounded_vector usedTracklets(mTimeFrame->getFoundTracklets(pivotRofId, 0).size(), 0, mMemoryPool.get()); trackletSelectionKernelHost( mTimeFrame->getClusters()[0].data(), mTimeFrame->getClusters()[1].data(), @@ -313,7 +318,8 @@ void VertexerTraits::computeTrackletMatching(const int iteration) mTimeFrame->getLinesLabel(pivotRofId), static_cast(mTimeFrame->getClustersOnLayer(pivotRofId, 1).size()), mVrtParams[iteration].tanLambdaCut, - mVrtParams[iteration].phiCut); + mVrtParams[iteration].phiCut, + constants::MaxSelectedTrackletsPerCluster); totalLines.local() += mTimeFrame->getLines(pivotRofId).size(); } }); From 3431463e58c0fb9aa79baa74ea9fd0cf669db07c Mon Sep 17 00:00:00 2001 From: Felix Schlepper Date: Sat, 25 Apr 2026 17:00:30 +0200 Subject: [PATCH 3/6] ITS: improve logging Signed-off-by: Felix Schlepper --- .../include/ITStracking/Definitions.h | 15 ++++- .../tracking/include/ITStracking/Tracker.h | 17 +++-- .../tracking/include/ITStracking/Vertexer.h | 26 ++++---- Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx | 33 ++++++++-- .../ITS/tracking/src/TrackingInterface.cxx | 1 + .../ITSMFT/ITS/tracking/src/Vertexer.cxx | 66 ++++++++++++++----- .../ITSMFT/ITS/workflow/src/TrackerSpec.cxx | 7 +- 7 files changed, 121 insertions(+), 44 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Definitions.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Definitions.h index 8dadf826aa80a..d79ea8c8bece8 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Definitions.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Definitions.h @@ -50,6 +50,19 @@ struct LogLogThrottler { return false; } }; + +struct TimingStats { + std::uint64_t calls = 0; + double totalTimeMs = 0.; + + void add(double timeMs) + { + ++calls; + totalTimeMs += timeMs; + } + double averageTimeMs() const { return calls ? totalTimeMs / static_cast(calls) : 0.; } +}; + } // namespace o2::its -#endif \ No newline at end of file +#endif diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h index ad8ea5b3b56af..240b0eb1e2f63 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h @@ -25,10 +25,12 @@ #include #include #include +#include #include #include "ITStracking/Configuration.h" +#include "ITStracking/Definitions.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/TrackerTraits.h" #include "ITStracking/BoundedAllocator.h" @@ -91,16 +93,18 @@ class Tracker double mTotalTime{0}; std::shared_ptr mMemoryPool; - enum State { + enum Steps { TFInit = 0, Trackleting, Celling, Neighbouring, Roading, - NStates, + NSteps, }; - State mCurState{TFInit}; - static constexpr std::array StateNames{"TimeFrame initialisation", "Tracklet finding", "Cell finding", "Neighbour finding", "Road finding"}; + Steps mCurStep{TFInit}; + static constexpr std::array StateNames{"TimeFrame initialisation", "Tracklet finding", "Cell finding", "Neighbour finding", "Road finding"}; + std::vector> mTimingStats; + void addTimingStatCurStep(int iteration, double timeMs); }; template @@ -125,7 +129,7 @@ float Tracker::evaluateTask(void (Tracker::*task)(T...), std:: } logger(sstream.str()); - if (mTrkParams[0].SaveTimeBenchmarks) { + if (mTrkParams[iteration].SaveTimeBenchmarks) { std::string taskNameStr(taskName); std::transform(taskNameStr.begin(), taskNameStr.end(), taskNameStr.begin(), [](unsigned char c) { return std::tolower(c); }); @@ -133,6 +137,7 @@ float Tracker::evaluateTask(void (Tracker::*task)(T...), std:: if (std::ofstream file{"its_time_benchmarks.txt", std::ios::app}) { file << "trk:" << iteration << '\t' << taskNameStr << '\t' << diff << '\n'; } + addTimingStatCurStep(iteration, diff); } } else { @@ -140,7 +145,7 @@ float Tracker::evaluateTask(void (Tracker::*task)(T...), std:: } if (mTrkParams[iteration].PrintMemory) { - LOGP(info, "iter:{}:{}: {}", iteration, StateNames[mCurState], mMemoryPool->asString()); + LOGP(info, "iter:{}:{}: {}", iteration, StateNames[mCurStep], mMemoryPool->asString()); } return diff; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h index f1cf081473264..eff91e820c56d 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h @@ -22,11 +22,11 @@ #include #include #include +#include #include #include "ITStracking/Constants.h" -#include "ITStracking/Definitions.h" #include "ITStracking/Configuration.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/VertexerTraits.h" @@ -56,6 +56,7 @@ class Vertexer float clustersToVertices(LogFunc = [](const std::string& s) { std::cout << s << '\n'; }); void filterMCTracklets(); + void printSummary() const; template void findTracklets(T&&... args) @@ -90,9 +91,9 @@ class Vertexer float evaluateTask(void (Vertexer::*task)(T...), std::string_view taskName, int iteration, LogFunc& logger, T&&... args); void printEpilog(LogFunc& logger, - const unsigned int trackletN01, const unsigned int trackletN12, - const unsigned selectedN, const unsigned int vertexN, const unsigned int totalVertexN, - const float trackletT, const float selecT, const float vertexT); + unsigned int trackletN01, unsigned int trackletN12, + unsigned selectedN, unsigned int vertexN, unsigned int totalVertexN, + float initT, float trackletT, float selecT, float vertexT); void setNThreads(int n, std::shared_ptr& arena) { mTraits->setNThreads(n, arena); } @@ -105,16 +106,18 @@ class Vertexer std::vector mVertParams; std::shared_ptr mMemoryPool; - enum State { + enum Steps { Init = 0, Trackleting, - Validating, + Selection, Finding, TruthSeeding, - NStates, + NSteps, }; - State mCurState{Init}; - static constexpr std::array StateNames{"Initialisation", "Tracklet finding", "Tracklet validation", "Vertex finding", "Truth seeding"}; + Steps mCurStep{Init}; + static constexpr std::array StateNames{"Initialisation", "Tracklet finding", "Tracklet selection", "Vertex finding", "Truth seeding"}; + std::vector> mTimingStats; + void addTimingStatCurStep(int iteration, double timeMs); }; template @@ -139,7 +142,7 @@ float Vertexer::evaluateTask(void (Vertexer::*task)(T...), std } logger(sstream.str()); - if (mVertParams[0].SaveTimeBenchmarks) { + if (mVertParams[iteration].SaveTimeBenchmarks) { std::string taskNameStr(taskName); std::transform(taskNameStr.begin(), taskNameStr.end(), taskNameStr.begin(), [](unsigned char c) { return std::tolower(c); }); @@ -147,13 +150,14 @@ float Vertexer::evaluateTask(void (Vertexer::*task)(T...), std if (std::ofstream file{"its_time_benchmarks.txt", std::ios::app}) { file << "vtx:" << iteration << '\t' << taskNameStr << '\t' << diff << '\n'; } + addTimingStatCurStep(iteration, diff); } } else { (this->*task)(std::forward(args)...); } if (mVertParams[iteration].PrintMemory) { - LOGP(info, "iter:{}:{}: {}", iteration, StateNames[mCurState], mMemoryPool->asString()); + LOGP(info, "iter:{}:{}: {}", iteration, StateNames[mCurStep], mMemoryPool->asString()); } return diff; diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index 3e91788c9881c..382f2314b2e6a 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -49,7 +49,7 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e int iteration{0}, iVertex{0}; auto handleException = [&](const auto& err) { LOGP(error, "Too much memory in {} in iteration {} iVtx={}: {:.2f} GB. Current limit is {:.2f} GB, check the detector status and/or the selections.", - StateNames[mCurState], iteration, iVertex, + StateNames[mCurStep], iteration, iVertex, (double)mTimeFrame->getArtefactsMemory() / GB, (double)mTrkParams[iteration].MaxMemory / GB); if (typeid(err) != typeid(std::bad_alloc)) { // only print if the exceptions is different from what is expected @@ -76,16 +76,16 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e int nTracks{-static_cast(mTimeFrame->getNumberOfTracks())}; iVertex = std::min(maxNvertices, 0); logger(std::format("==== ITS {} Tracking iteration {} summary ====", mTraits->getName(), iteration)); - total += timeFrame = evaluateTask(&Tracker::initialiseTimeFrame, StateNames[mCurState = TFInit], iteration, evalLog, iteration); + total += timeFrame = evaluateTask(&Tracker::initialiseTimeFrame, StateNames[mCurStep = TFInit], iteration, evalLog, iteration); logger(std::format(" - TimeFrame initialisation completed in {:.2f} ms", timeFrame)); do { - timeTracklets += evaluateTask(&Tracker::computeTracklets, StateNames[mCurState = Trackleting], iteration, evalLog, iteration, iVertex); + timeTracklets += evaluateTask(&Tracker::computeTracklets, StateNames[mCurStep = Trackleting], iteration, evalLog, iteration, iVertex); nTracklets += mTraits->getTFNumberOfTracklets(); - timeCells += evaluateTask(&Tracker::computeCells, StateNames[mCurState = Celling], iteration, evalLog, iteration); + timeCells += evaluateTask(&Tracker::computeCells, StateNames[mCurStep = Celling], iteration, evalLog, iteration); nCells += mTraits->getTFNumberOfCells(); - timeNeighbours += evaluateTask(&Tracker::findCellsNeighbours, StateNames[mCurState = Neighbouring], iteration, evalLog, iteration); + timeNeighbours += evaluateTask(&Tracker::findCellsNeighbours, StateNames[mCurStep = Neighbouring], iteration, evalLog, iteration); nNeighbours += mTimeFrame->getNumberOfNeighbours(); - timeRoads += evaluateTask(&Tracker::findRoads, StateNames[mCurState = Roading], iteration, evalLog, iteration); + timeRoads += evaluateTask(&Tracker::findRoads, StateNames[mCurStep = Roading], iteration, evalLog, iteration); } while (++iVertex < maxNvertices); logger(std::format(" - Tracklet finding: {} tracklets found in {:.2f} ms", nTracklets, timeTracklets)); logger(std::format(" - Cell finding: {} cells found in {:.2f} ms", nCells, timeCells)); @@ -229,12 +229,33 @@ void Tracker::adoptTimeFrame(TimeFrame& tf) mTraits->adoptTimeFrame(&tf); } +template +void Tracker::addTimingStatCurStep(int iteration, double timeMs) +{ + if (iteration < 0) { + return; + } + if (mTimingStats.size() < (iteration + 1)) { + mTimingStats.resize(iteration + 1); + } + mTimingStats[iteration][mCurStep].add(timeMs); +} + template void Tracker::printSummary() const { auto avgTF = mTotalTime * 1.e-3 / ((mTimeFrameCounter > 0) ? (double)mTimeFrameCounter : -1.0); auto avgTFwithDropped = mTotalTime * 1.e-3 / (((mTimeFrameCounter + mNumberOfDroppedTFs) > 0) ? (double)(mTimeFrameCounter + mNumberOfDroppedTFs) : -1.0); LOGP(info, "Tracker summary: Processed {} TFs (dropped {}) in TOT={:.2f} s, AVG/TF={:.2f} ({:.2f}) s", mTimeFrameCounter, mNumberOfDroppedTFs, mTotalTime * 1.e-3, avgTF, avgTFwithDropped); + for (size_t iteration = 0; iteration < mTimingStats.size(); ++iteration) { + for (size_t state = 0; state < NSteps; ++state) { + const auto& stats = mTimingStats[iteration][state]; + if (!stats.calls) { + continue; + } + LOGP(info, " - iter {} {}: calls={} total={:.2f} ms avg={:.2f} ms", iteration, StateNames[state], stats.calls, stats.totalTimeMs, stats.averageTimeMs()); + } + } } template class Tracker<7>; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx index cc8731c8b6912..af82e326536a8 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx @@ -457,6 +457,7 @@ void ITSTrackingInterface::finaliseCCDB(ConcreteDataMatcher& matcher, void* obj) void ITSTrackingInterface::printSummary() const { + mVertexer->printSummary(); mTracker->printSummary(); } diff --git a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx index 556302cb2854f..2acbec7fe8bc0 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx @@ -15,11 +15,6 @@ #include "ITStracking/Vertexer.h" #include "ITStracking/BoundedAllocator.h" -#include "ITStracking/Cluster.h" - -#include "ITStracking/ClusterLines.h" -#include "ITStracking/Tracklet.h" -#include "ITStracking/IndexTableUtils.h" #include "ITStracking/VertexerTraits.h" #include "ITStracking/TrackingConfigParam.h" @@ -41,8 +36,9 @@ float Vertexer::clustersToVertices(LogFunc logger) LogFunc evalLog = [](const std::string&) {}; if (mTimeFrame->hasMCinformation() && mVertParams[0].useTruthSeeding) { - float t = evaluateTask(&Vertexer::addTruthSeeds, StateNames[mCurState = TruthSeeding], 0, evalLog); + float t = evaluateTask(&Vertexer::addTruthSeeds, StateNames[mCurStep = TruthSeeding], 0, evalLog); sortVertices(); + ++mTimeFrameCounter; return t; } @@ -50,7 +46,7 @@ float Vertexer::clustersToVertices(LogFunc logger) mTraits->updateVertexingParameters(mVertParams); auto handleException = [&](const auto& err) { - LOGP(error, "Encountered critical error in step {}, stopping further processing of this TF: {}", StateNames[mCurState], err.what()); + LOGP(error, "Encountered critical error in step {}, stopping further processing of this TF: {}", StateNames[mCurStep], err.what()); if (!mVertParams[0].DropTFUponFailure) { throw err; } else { @@ -59,6 +55,7 @@ float Vertexer::clustersToVertices(LogFunc logger) }; float timeTracklet{0.f}, timeSelection{0.f}, timeVertexing{0.f}, timeInit{0.f}; + bool completed = false; try { for (int iteration = 0; iteration < (int)mVertParams.size(); ++iteration) { mMemoryPool->setMaxMemory(mVertParams[iteration].MaxMemory); @@ -66,15 +63,15 @@ float Vertexer::clustersToVertices(LogFunc logger) logger(fmt::format("=== ITS {} Seeding vertexer iteration {} summary:", mTraits->getName(), iteration)); trkPars.PhiBins = mTraits->getVertexingParameters()[0].PhiBins; trkPars.ZBins = mTraits->getVertexingParameters()[0].ZBins; - auto timeInitIteration = evaluateTask(&Vertexer::initialiseVertexer, StateNames[mCurState = Init], iteration, evalLog, trkPars, iteration); - auto timeTrackletIteration = evaluateTask(&Vertexer::findTracklets, StateNames[mCurState = Trackleting], iteration, evalLog, iteration); + auto timeInitIteration = evaluateTask(&Vertexer::initialiseVertexer, StateNames[mCurStep = Init], iteration, evalLog, trkPars, iteration); + auto timeTrackletIteration = evaluateTask(&Vertexer::findTracklets, StateNames[mCurStep = Trackleting], iteration, evalLog, iteration); nTracklets01 = mTimeFrame->getTotalTrackletsTF(0); nTracklets12 = mTimeFrame->getTotalTrackletsTF(1); - auto timeSelectionIteration = evaluateTask(&Vertexer::validateTracklets, StateNames[mCurState = Validating], iteration, evalLog, iteration); + auto timeSelectionIteration = evaluateTask(&Vertexer::validateTracklets, StateNames[mCurStep = Selection], iteration, evalLog, iteration); const auto nVerticesBefore = mTimeFrame->getPrimaryVertices().size(); - auto timeVertexingIteration = evaluateTask(&Vertexer::findVertices, StateNames[mCurState = Finding], iteration, evalLog, iteration); + auto timeVertexingIteration = evaluateTask(&Vertexer::findVertices, StateNames[mCurStep = Finding], iteration, evalLog, iteration); const auto nVerticesAfter = mTimeFrame->getPrimaryVertices().size(); - printEpilog(logger, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), nVerticesAfter - nVerticesBefore, nVerticesAfter, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); + printEpilog(logger, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), nVerticesAfter - nVerticesBefore, nVerticesAfter, timeInitIteration, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); timeInit += timeInitIteration; timeTracklet += timeTrackletIteration; timeSelection += timeSelectionIteration; @@ -84,6 +81,7 @@ float Vertexer::clustersToVertices(LogFunc logger) sortVertices(); mTimeFrame->updateROFVertexLookupTable(); } + completed = true; } catch (const BoundedMemoryResource::MemoryLimitExceeded& err) { handleException(err); } catch (const std::bad_alloc& err) { @@ -92,6 +90,10 @@ float Vertexer::clustersToVertices(LogFunc logger) LOGP(fatal, "Uncaught exception!"); } + if (completed) { + ++mTimeFrameCounter; + } + return timeInit + timeTracklet + timeSelection + timeVertexing; } @@ -135,15 +137,43 @@ void Vertexer::adoptTimeFrame(TimeFrameN& tf) mTraits->adoptTimeFrame(&tf); } +template +void Vertexer::addTimingStatCurStep(int iteration, double timeMs) +{ + if (iteration < 0) { + return; + } + if (mTimingStats.size() < (iteration + 1)) { + mTimingStats.resize(iteration + 1); + } + mTimingStats[iteration][mCurStep].add(timeMs); +} + +template +void Vertexer::printSummary() const +{ + LOGP(info, "Vertexer summary: Processed {} TFs", mTimeFrameCounter); + for (size_t iteration = 0; iteration < mTimingStats.size(); ++iteration) { + for (size_t state = 0; state < NSteps; ++state) { + const auto& stats = mTimingStats[iteration][state]; + if (!stats.calls) { + continue; + } + LOGP(info, " - iter {} {}: calls={} total={:.2f} ms avg={:.2f} ms", iteration, StateNames[state], stats.calls, stats.totalTimeMs, stats.averageTimeMs()); + } + } +} + template void Vertexer::printEpilog(LogFunc& logger, - const unsigned int trackletN01, const unsigned int trackletN12, - const unsigned selectedN, const unsigned int vertexN, const unsigned int totalVertexN, - const float trackletT, const float selecT, const float vertexT) + unsigned int trackletN01, unsigned int trackletN12, + unsigned selectedN, unsigned int vertexN, unsigned int totalVertexN, + float initT, float trackletT, float selecT, float vertexT) { - logger(fmt::format(" - {} Vertexer: found {} | {} tracklets in: {:.2f} ms", mTraits->getName(), trackletN01, trackletN12, trackletT)); - logger(fmt::format(" - {} Vertexer: selected {} tracklets in: {:.2f} ms", mTraits->getName(), selectedN, selecT)); - logger(fmt::format(" - {} Vertexer: found {} vertices (total: {}) in: {:.2f} ms", mTraits->getName(), vertexN, totalVertexN, vertexT)); + logger(fmt::format(" - {}: completed in {:.2f} ms", StateNames[Init], initT)); + logger(fmt::format(" - {}: found {} | {} tracklets in {:.2f} ms", StateNames[Trackleting], trackletN01, trackletN12, trackletT)); + logger(fmt::format(" - {}: selected {} tracklets in {:.2f} ms", StateNames[Selection], selectedN, selecT)); + logger(fmt::format(" - {}: found {} vertices (total {}) in {:.2f} ms", StateNames[Finding], vertexN, totalVertexN, vertexT)); } template class Vertexer<7>; diff --git a/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx b/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx index 932c82c2d1ca4..bbafc48e931ed 100644 --- a/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx +++ b/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx @@ -85,8 +85,11 @@ void TrackerDPL::endOfStream(EndOfStreamContext& ec) void TrackerDPL::end() { - mITSTrackingInterface.printSummary(); - LOGF(info, "ITS CA-Tracker total timing: Cpu: %.3e Real: %.3e s in %d slots", mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1); + if (static bool printOnce{false}; !printOnce) { + printOnce = true; + mITSTrackingInterface.printSummary(); + LOGF(info, "ITS CA-Tracker total timing: Cpu: %.3e Real: %.3e s in %d slots", mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1); + } } DataProcessorSpec getTrackerSpec(bool useMC, bool doStag, bool useGeom, int trgType, TrackingMode::Type trMode, const bool overrBeamEst, o2::gpu::gpudatatypes::DeviceType dType) From 39f2fbf18fe3cc860bc106b9b621f4a4da5f4e02 Mon Sep 17 00:00:00 2001 From: Felix Schlepper Date: Sat, 25 Apr 2026 20:31:04 +0200 Subject: [PATCH 4/6] ITS: cleanup tracklet class Signed-off-by: Felix Schlepper --- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 18 +---- .../ITS/tracking/include/ITStracking/Cell.h | 4 +- .../tracking/include/ITStracking/Constants.h | 1 + .../tracking/include/ITStracking/Tracklet.h | 69 +++++++------------ .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 16 ++--- .../ITS/tracking/src/VertexerTraits.cxx | 20 +++--- 6 files changed, 41 insertions(+), 87 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 7cf0ef4c0b0d7..49b8f19d68ea6 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -47,20 +47,6 @@ namespace o2::its namespace gpu { -struct sort_tracklets { - GPUhd() bool operator()(const Tracklet& a, const Tracklet& b) - { - if (a.firstClusterIndex != b.firstClusterIndex) { - return a.firstClusterIndex < b.firstClusterIndex; - } - return a.secondClusterIndex < b.secondClusterIndex; - } -}; - -struct equal_tracklets { - GPUhd() bool operator()(const Tracklet& a, const Tracklet& b) { return a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex == b.secondClusterIndex; } -}; - template struct sort_by_second { GPUhd() bool operator()(const gpuPair& a, const gpuPair& b) const { return a.second < b.second; } @@ -687,8 +673,8 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, mulScatAng[layer]); thrust::device_ptr tracklets_ptr(spanTracklets[layer]); auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator(alloc)).on(streams[layer].get()); - thrust::sort(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer], gpu::sort_tracklets()); - auto unique_end = thrust::unique(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer], gpu::equal_tracklets()); + thrust::sort(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer]); + auto unique_end = thrust::unique(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[layer]); nTracklets[layer] = unique_end - tracklets_ptr; if (layer) { GPUChkErrS(cudaMemsetAsync(trackletsLUTsHost[layer], 0, (nClusters[layer] + 1) * sizeof(int), streams[layer].get())); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h index 0cd81082eda32..c7718ee666311 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h @@ -56,8 +56,8 @@ class SeedBase : public o2::track::TrackParCovF GPUhd() const auto& clustersRaw() const { return mClusters; } private: - float mChi2 = -999.f; - int mLevel = constants::UnusedIndex; + float mChi2{constants::UnsetValue}; + int mLevel{constants::UnusedIndex}; std::array mTracklets = constants::helpers::initArray(); std::array mClusters = constants::helpers::initArray(); TimeEstBC mTime; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h index 85f4b3a222af2..34fa819b178eb 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h @@ -30,6 +30,7 @@ constexpr bool SaveTimeBenchmarks = false; constexpr float Tolerance = 1e-12; // numerical tolerance constexpr int ClustersPerCell = 3; // number of clusters for a cell 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] constexpr float Rho = 2.33f; // Density of Si [g/cm^3] constexpr int MaxIter = 4; // Max. supported iterations diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h index d93a5e1c7d70e..829fe9fa984e4 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h @@ -24,27 +24,35 @@ #include "GPUCommonDef.h" #include "GPUCommonLogger.h" -#ifndef GPUCA_GPUCODE_DEVICE -#ifndef GPU_NO_FMT -#include -#include -#endif -#endif - namespace o2::its { +// tracklets are entirely determined by their two cluster idx struct Tracklet final { GPUhdDefault() Tracklet() = default; - GPUhdi() Tracklet(const int, const int, const Cluster&, const Cluster&, const TimeEstBC& t); - GPUhdi() Tracklet(const int, const int, float tanL, float phi, const TimeEstBC& t); - GPUhdDefault() bool operator==(const Tracklet&) const = default; - GPUhdi() unsigned char isEmpty() const + GPUhdi() Tracklet(const int firstClusterOrderingIndex, const int secondClusterOrderingIndex, + const Cluster& firstCluster, const Cluster& secondCluster, const TimeEstBC& t) + : firstClusterIndex(firstClusterOrderingIndex), + secondClusterIndex(secondClusterOrderingIndex), + tanLambda((firstCluster.zCoordinate - secondCluster.zCoordinate) / (firstCluster.radius - secondCluster.radius)), + phi(o2::gpu::GPUCommonMath::ATan2(firstCluster.yCoordinate - secondCluster.yCoordinate, firstCluster.xCoordinate - secondCluster.xCoordinate)), + mTime(t) {} + + GPUhdi() Tracklet(const int idx0, const int idx1, float tanL, float phi, const TimeEstBC& t) + : firstClusterIndex(idx0), + secondClusterIndex(idx1), + tanLambda(tanL), + phi(phi), + mTime(t) {} + GPUhdi() bool operator<(const Tracklet& o) const noexcept + { + return (firstClusterIndex != o.firstClusterIndex) ? firstClusterIndex < o.firstClusterIndex : secondClusterIndex < o.secondClusterIndex; + } + GPUhdi() bool operator==(const Tracklet& o) const noexcept { - return firstClusterIndex < 0 || secondClusterIndex < 0; + return firstClusterIndex == o.firstClusterIndex && secondClusterIndex == o.secondClusterIndex; } GPUhdi() bool isCompatible(const Tracklet& o) const { return mTime.isCompatible(o.mTime); } - GPUhdi() unsigned char operator<(const Tracklet&) const; GPUhd() void print() const { LOGP(info, "TRKLT: fClIdx:{} sClIdx:{} ts:{}+/-{} TgL={} Phi={}", firstClusterIndex, secondClusterIndex, mTime.getTimeStamp(), mTime.getTimeStampError(), tanLambda, phi); @@ -54,44 +62,13 @@ struct Tracklet final { int firstClusterIndex{constants::UnusedIndex}; int secondClusterIndex{constants::UnusedIndex}; - float tanLambda{-999}; - float phi{-999}; + float tanLambda{constants::UnsetValue}; + float phi{constants::UnsetValue}; TimeEstBC mTime; ClassDefNV(Tracklet, 1); }; -GPUhdi() Tracklet::Tracklet(const int firstClusterOrderingIndex, const int secondClusterOrderingIndex, - const Cluster& firstCluster, const Cluster& secondCluster, const TimeEstBC& t) - : firstClusterIndex(firstClusterOrderingIndex), - secondClusterIndex(secondClusterOrderingIndex), - tanLambda((firstCluster.zCoordinate - secondCluster.zCoordinate) / - (firstCluster.radius - secondCluster.radius)), - phi(o2::gpu::GPUCommonMath::ATan2(firstCluster.yCoordinate - secondCluster.yCoordinate, - firstCluster.xCoordinate - secondCluster.xCoordinate)), - mTime(t) -{ - // Nothing to do -} - -GPUhdi() Tracklet::Tracklet(const int idx0, const int idx1, float tanL, float phi, const TimeEstBC& t) - : firstClusterIndex(idx0), - secondClusterIndex(idx1), - tanLambda(tanL), - phi(phi), - mTime(t) -{ - // Nothing to do -} - -GPUhdi() unsigned char Tracklet::operator<(const Tracklet& t) const -{ - if (isEmpty()) { - return false; - } - return true; -} - } // namespace o2::its #endif /* TRACKINGITS_INCLUDE_TRACKLET_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index d5db4b983f754..337fffcdcc87e 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -206,19 +206,11 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer } tbb::parallel_for(0, mTrkParams[iteration].TrackletsPerRoad(), [&](const int iLayer) { - /// Sort tracklets + /// Sort tracklets & remove duplicates + // duplicates can exist simply since we evaluate per vertex auto& trkl{mTimeFrame->getTracklets()[iLayer]}; - std::sort(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) -> bool { - if (a.firstClusterIndex != b.firstClusterIndex) { - return a.firstClusterIndex < b.firstClusterIndex; - } - return a.secondClusterIndex < b.secondClusterIndex; - }); - /// Remove duplicates - trkl.erase(std::unique(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) -> bool { - return a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex == b.secondClusterIndex; - }), - trkl.end()); + std::sort(trkl.begin(), trkl.end()); + trkl.erase(std::unique(trkl.begin(), trkl.end()), trkl.end()); trkl.shrink_to_fit(); if (iLayer > 0) { /// recalculate lut auto& lut{mTimeFrame->getTrackletsLookupTable()[iLayer - 1]}; diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index 6e0b406cddfd1..00674b715b97d 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -269,20 +269,18 @@ void VertexerTraits::computeTracklets(const int iteration) if (mTimeFrame->hasMCinformation()) { for (const auto& trk : mTimeFrame->getTracklets()[0]) { o2::MCCompLabel label; - if (!trk.isEmpty()) { - int sortedId0{trk.firstClusterIndex}; - int sortedId1{trk.secondClusterIndex}; - for (const auto& lab0 : mTimeFrame->getClusterLabels(0, mTimeFrame->getClusters()[0][sortedId0].clusterId)) { - for (const auto& lab1 : mTimeFrame->getClusterLabels(1, mTimeFrame->getClusters()[1][sortedId1].clusterId)) { - if (lab0 == lab1 && lab0.isValid()) { - label = lab0; - break; - } - } - if (label.isValid()) { + int sortedId0{trk.firstClusterIndex}; + int sortedId1{trk.secondClusterIndex}; + for (const auto& lab0 : mTimeFrame->getClusterLabels(0, mTimeFrame->getClusters()[0][sortedId0].clusterId)) { + for (const auto& lab1 : mTimeFrame->getClusterLabels(1, mTimeFrame->getClusters()[1][sortedId1].clusterId)) { + if (lab0 == lab1 && lab0.isValid()) { + label = lab0; break; } } + if (label.isValid()) { + break; + } } mTimeFrame->getTrackletsLabel(0).emplace_back(label); } From de234f4d15f0c1112a92b20cd28fdae2e1b45ba7 Mon Sep 17 00:00:00 2001 From: Felix Schlepper Date: Sun, 26 Apr 2026 19:36:14 +0200 Subject: [PATCH 5/6] ITS: let staggering macro also run over non-staggered data Signed-off-by: Felix Schlepper --- .../ITSMFT/ITS/macros/test/CheckStaggering.C | 18 ++++++++++++------ 1 file changed, 12 insertions(+), 6 deletions(-) diff --git a/Detectors/ITSMFT/ITS/macros/test/CheckStaggering.C b/Detectors/ITSMFT/ITS/macros/test/CheckStaggering.C index e3a79779a5fb1..11ebcbbb1b1f1 100644 --- a/Detectors/ITSMFT/ITS/macros/test/CheckStaggering.C +++ b/Detectors/ITSMFT/ITS/macros/test/CheckStaggering.C @@ -80,9 +80,9 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "") auto& ccdbmgr = o2::ccdb::BasicCCDBManager::instance(); ccdbmgr.setURL("https://alice-ccdb.cern.ch"); auto runDuration = ccdbmgr.getRunDuration(runNumber); - auto tRun = runDuration.first + (runDuration.second - runDuration.first) / 2; // time stamp for the middle of the run duration + auto tRun = runDuration.first + ((runDuration.second - runDuration.first) / 2); // time stamp for the middle of the run duration ccdbmgr.setTimestamp(tRun); - printf("Run %d has TS %lld", runNumber, tRun); + printf("Run %d has TS %ld", runNumber, tRun); auto geoAligned = ccdbmgr.get("GLO/Config/GeometryAligned"); auto magField = ccdbmgr.get("GLO/Config/GRPMagField"); auto grpLHC = ccdbmgr.get("GLO/Config/GRPLHCIF"); @@ -134,7 +134,7 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "") auto hVtxZ = new TH1F("hVtxZ", "seeding vertices Z", 200, -16, 16); auto hVtxNCont = new TH1F("hVtxNCont", "seeding vertices contributors", 100, 0, 100); auto hVtxZNCont = new TProfile("hVtxZNCont", "seeding vertices z-contributors", 200, -16, 16); - auto hVtxCls = new TProfile("hVtxCls", ";Cls/TF;Cls/Vtx", 400, 20000, 60000); + auto hVtxCls = new TProfile("hVtxCls", ";Cls/TF;Cls/Vtx", 2000, 600000, 900000); auto hVtxTS = new TH1D("hVtxTS", "vtx time t0;t0 (BC)", o2::constants::lhc::LHCMaxBunches, 0, o2::constants::lhc::LHCMaxBunches); const float minVtxWeight{5}; @@ -183,8 +183,12 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "") tTrks->SetBranchAddress("ITSTrack", &trkArrPtr); tTrks->SetBranchAddress("Vertices", &vtxArrPtr); - for (int i{0}; i < 7; ++i) { - tCls->SetBranchAddress(Form("ITSClusterComp_%d", i), &clsArr[i]); + if (tCls->GetBranchStatus("ITSClusterComp")) { + tCls->SetBranchAddress("ITSClusterComp", &clsArr[0]); + } else { + for (int i{0}; i < 7; ++i) { + tCls->SetBranchAddress(Form("ITSClusterComp_%d", i), &clsArr[i]); + } } for (int iTF{0}; tTrks->LoadTree(iTF) >= 0; ++iTF) { @@ -193,7 +197,9 @@ void CheckStaggering(int runNumber, int max = -1, const std::string& dir = "") size_t ncls = 0; for (int i{0}; i < 7; ++i) { - ncls += clsArr[i]->size(); + if (clsArr[i]) { + ncls += clsArr[i]->size(); + } } // for each TF built pool of positive and negaitve tracks From db833d6287b2b9b5571b5a57ebdd2d9b15d0025a Mon Sep 17 00:00:00 2001 From: Felix Schlepper Date: Sun, 26 Apr 2026 21:19:47 +0200 Subject: [PATCH 6/6] ITS: fix time assignment Signed-off-by: Felix Schlepper --- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 28 +++++--- .../ITS/tracking/test/testROFLookupTables.cxx | 72 +++++++++++++++++++ 2 files changed, 91 insertions(+), 9 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 337fffcdcc87e..dc2d6e8889973 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -801,30 +801,40 @@ void TrackerTraits::acceptTracks(int iteration, bounded_vectormarkUsedCluster(iLayer, track.getClusterIndex(iLayer)); int currentROF = mTimeFrame->getClusterROF(iLayer, track.getClusterIndex(iLayer)); - auto rofTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(currentROF, true); + const auto nominalROFTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(currentROF); + const auto expandedROFTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(currentROF, true); if (firstCls) { firstCls = false; - ts = rofTS; + nominalTS = nominalROFTS; + expandedTS = expandedROFTS; } else { - if (!ts.isCompatible(rofTS)) { - LOGP(fatal, "TS {}+/-{} are incompatible with {}+/-{}, this should not happen!", rofTS.getTimeStamp(), rofTS.getTimeStampError(), ts.getTimeStamp(), ts.getTimeStampError()); + if (nominalCompatible) { + if (nominalTS.isCompatible(nominalROFTS)) { + nominalTS += nominalROFTS; + } else { + nominalCompatible = false; + } } - ts += rofTS; + if (!expandedTS.isCompatible(expandedROFTS)) { + LOGP(fatal, "TS {}+/-{} are incompatible with {}+/-{}, this should not happen!", expandedROFTS.getTimeStamp(), expandedROFTS.getTimeStampError(), expandedTS.getTimeStamp(), expandedTS.getTimeStampError()); + } + expandedTS += expandedROFTS; } } - track.getTimeStamp() = ts.makeSymmetrical(); + track.getTimeStamp() = (nominalCompatible ? nominalTS : expandedTS).makeSymmetrical(); + // this is a sanity clamp + // we cannot be worse than the clock so we clamp to this if (track.getTimeStamp().getTimeStampError() > smallestROFHalf) { track.getTimeStamp().setTimeStampError(smallestROFHalf); } - track.setUserField(0); track.getParamOut().setUserField(0); mTimeFrame->getTracks().emplace_back(track); diff --git a/Detectors/ITSMFT/ITS/tracking/test/testROFLookupTables.cxx b/Detectors/ITSMFT/ITS/tracking/test/testROFLookupTables.cxx index 8594e59149444..dd98a75efca7c 100644 --- a/Detectors/ITSMFT/ITS/tracking/test/testROFLookupTables.cxx +++ b/Detectors/ITSMFT/ITS/tracking/test/testROFLookupTables.cxx @@ -135,6 +135,78 @@ BOOST_AUTO_TEST_CASE(rofoverlap_staggered_pp) view.printAll(); } +BOOST_AUTO_TEST_CASE(rofoverlap_staggered_track_time_ignores_added_error) +{ + const uint32_t rofLen{198}; + const uint32_t rofDelay{33}; + const uint32_t addTimeErr{100}; + + o2::its::ROFOverlapTable<7> tableNoError; + o2::its::ROFOverlapTable<7> tableWithError; + for (uint32_t lay{0}; lay < 7; ++lay) { + const auto delay = (lay == 6) ? 0 : lay * rofDelay; + tableNoError.defineLayer(lay, 2, rofLen, delay, 0, 0); + tableWithError.defineLayer(lay, 2, rofLen, delay, 0, addTimeErr); + } + + auto getCommonTrackTime = [](const auto& table) { + auto ts = table.getLayer(0).getROFTimeBounds(0); + for (uint32_t lay{1}; lay < 7; ++lay) { + ts += table.getLayer(lay).getROFTimeBounds(0); + } + return ts.makeSymmetrical(); + }; + + const auto tsNoError = getCommonTrackTime(tableNoError); + BOOST_CHECK_EQUAL(tsNoError.getTimeStamp(), 181.5f); + BOOST_CHECK_EQUAL(tsNoError.getTimeStampError(), 16.5f); + + const auto tsWithError = getCommonTrackTime(tableWithError); + BOOST_CHECK_EQUAL(tsWithError.getTimeStamp(), 181.5f); + BOOST_CHECK_EQUAL(tsWithError.getTimeStampError(), 16.5f); +} + +BOOST_AUTO_TEST_CASE(rofoverlap_track_time_boundary_migration_fallback) +{ + const uint32_t rofLen{198}; + const uint32_t addTimeErr{30}; + + o2::its::ROFOverlapTable<7> table; + for (uint32_t lay{0}; lay < 7; ++lay) { + table.defineLayer(lay, 4, rofLen, 0, 0, addTimeErr); + } + + auto getCommonTrackTime = [](const auto& table) { + bool firstCls{true}, nominalCompatible{true}; + o2::its::TimeEstBC nominalTS, expandedTS; + for (uint32_t lay{0}; lay < 7; ++lay) { + const auto rof = lay < 3 ? 0 : 1; + const auto nominalROFTS = table.getLayer(lay).getROFTimeBounds(rof); + const auto expandedROFTS = table.getLayer(lay).getROFTimeBounds(rof, true); + if (firstCls) { + firstCls = false; + nominalTS = nominalROFTS; + expandedTS = expandedROFTS; + } else { + if (nominalCompatible) { + if (nominalTS.isCompatible(nominalROFTS)) { + nominalTS += nominalROFTS; + } else { + nominalCompatible = false; + } + } + BOOST_REQUIRE(expandedTS.isCompatible(expandedROFTS)); + expandedTS += expandedROFTS; + } + } + return (nominalCompatible ? nominalTS : expandedTS).makeSymmetrical(); + }; + + const auto tsWithError = getCommonTrackTime(table); + BOOST_CHECK_EQUAL(tsWithError.getTimeStamp(), 198.f); + BOOST_CHECK_EQUAL(tsWithError.getTimeStampError(), 30.f); +} + BOOST_AUTO_TEST_CASE(rofoverlap_staggered_alllayers) { // test staggered layers with ROF delay