From 4568c80076ef990b9555fc675845051c3941f17e Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Thu, 21 May 2026 14:03:39 +0200 Subject: [PATCH 1/4] ITS: re-enable the possibility of extending tracks --- .../GPU/ITStrackingGPU/TimeFrameGPU.h | 23 + .../GPU/ITStrackingGPU/TrackerTraitsGPU.h | 6 + .../GPU/ITStrackingGPU/TrackingKernels.h | 54 +++ .../ITS/tracking/GPU/cuda/TimeFrameGPU.cu | 78 ++++ .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 124 ++++- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 425 ++++++++++++++++++ .../include/ITStracking/Configuration.h | 7 +- .../include/ITStracking/ROFLookupTables.h | 43 +- .../tracking/include/ITStracking/TimeFrame.h | 1 + .../ITStracking/TrackExtensionCandidate.h | 114 +++++ .../include/ITStracking/TrackFollower.h | 301 +++++++++++++ .../tracking/include/ITStracking/Tracker.h | 4 +- .../include/ITStracking/TrackerTraits.h | 57 +++ .../include/ITStracking/TrackingConfigParam.h | 4 + .../ITSMFT/ITS/tracking/src/Configuration.cxx | 19 + .../ITSMFT/ITS/tracking/src/TimeFrame.cxx | 1 + Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx | 13 + .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 378 ++++++++++++++++ 18 files changed, 1646 insertions(+), 6 deletions(-) create mode 100644 Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h create mode 100644 Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index 5f56e3f272473..7223968c8cbf9 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -19,6 +19,7 @@ #include "ITStracking/BoundedAllocator.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/Configuration.h" +#include "ITStracking/TrackExtensionCandidate.h" #include "ITStrackingGPU/Utils.h" namespace o2::its::gpu @@ -90,8 +91,13 @@ class TimeFrameGPU : public TimeFrame void createNeighboursDevice(const unsigned int layer); void createNeighboursLUTDevice(const int, const unsigned int); void createTrackITSExtDevice(const size_t); + void loadTrackExtensionStartTracksDevice(); + void createTrackExtensionCandidatesDevice(const size_t); + void createTrackExtensionScratchDevice(const int nThreads, const int beamWidth); + void createTrackExtensionResultsDevice(const size_t); void downloadTrackITSExtDevice(); void downloadCellsNeighboursDevice(std::vector>&, const int); + void downloadTrackExtensionResultsDevice(); void downloadNeighboursLUTDevice(bounded_vector&, const int); void downloadCellsDevice(); void downloadCellsLUTDevice(); @@ -118,6 +124,7 @@ class TimeFrameGPU : public TimeFrame const auto getDeviceTrackingTopologyView() const { return mDeviceTrackingTopologyView; } int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; } auto& getTrackITSExt() { return mTrackITSExt; } + auto& getTrackExtensionResults() { return mTrackExtensionResults; } Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; } int* getDeviceROFramesPV() { return mROFramesPVDevice; } unsigned char* getDeviceUsedClusters(const int); @@ -125,6 +132,12 @@ class TimeFrameGPU : public TimeFrame // Hybrid TrackITSExt* getDeviceTrackITSExt() { return mTrackITSExtDevice; } + TrackITSExt* getDeviceTrackExtensionStartTracks() { return mTrackExtensionStartTracksDevice; } + TrackExtensionCandidate* getDeviceTrackExtensionCandidates() { return mTrackExtensionCandidatesDevice; } + int* getDeviceTrackExtensionCandidateOffsets() { return mTrackExtensionCandidateOffsetsDevice; } + TrackExtensionHypothesis* getDeviceActiveTrackExtensionHypotheses() { return mActiveTrackExtensionHypothesesDevice; } + TrackExtensionHypothesis* getDeviceNextTrackExtensionHypotheses() { return mNextTrackExtensionHypothesesDevice; } + TrackExtensionResult* getDeviceTrackExtensionResults() { return mTrackExtensionResultsDevice; } int* getDeviceNeighboursLUT(const int layer) { return mNeighboursLUTDevice[layer]; } gsl::span getDeviceNeighboursLUTs() { return mNeighboursLUTDevice; } CellNeighbour** getDeviceArrayNeighbours() { return mNeighboursDeviceArray; } @@ -222,6 +235,13 @@ class TimeFrameGPU : public TimeFrame float** mCellSeedsChi2DeviceArray; TrackITSExt* mTrackITSExtDevice; + TrackITSExt* mTrackExtensionStartTracksDevice{nullptr}; + TrackExtensionCandidate* mTrackExtensionCandidatesDevice{nullptr}; + int* mTrackExtensionCandidateOffsetsDevice{nullptr}; + TrackExtensionHypothesis* mActiveTrackExtensionHypothesesDevice{nullptr}; + TrackExtensionHypothesis* mNextTrackExtensionHypothesesDevice{nullptr}; + TrackExtensionResult* mTrackExtensionResultsDevice{nullptr}; + unsigned int mNTrackExtensionResults{0}; std::array mNeighboursDevice{}; CellNeighbour** mNeighboursDeviceArray{nullptr}; std::array mTrackingFrameInfoDevice; @@ -238,6 +258,9 @@ class TimeFrameGPU : public TimeFrame // Temporary buffer for storing output tracks from GPU tracking bounded_vector mTrackITSExt; + bounded_vector mTrackExtensionStartTracks; + // Temporary buffer for fitted track extension proposals from GPU tracking + bounded_vector> mTrackExtensionResults; }; template diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h index 81d870c5b46c2..13773ac234027 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h @@ -35,6 +35,7 @@ class TrackerTraitsGPU final : public TrackerTraits void computeLayerCells(const int iteration) final; void findCellsNeighbours(const int iteration) final; void findRoads(const int iteration) final; + void extendTracks(const int iteration) final; void setBz(float) final; @@ -47,6 +48,11 @@ class TrackerTraitsGPU final : public TrackerTraits int getTFNumberOfCells() const override; private: + bool hasTrackFollower(const int iteration) const; + + void buildTrackExtensionCandidates(const int iteration, typename TrackerTraits::TrackExtensionCandidates& candidatesPerTrack) final; + bool materializeTrackExtensionCandidate(TrackITSExt& track, const typename TrackerTraits::TrackExtensionCandidateN& candidate, const int iteration) final; + IndexTableUtilsN* mDeviceIndexTableUtils; gpu::TimeFrameGPU* mTimeFrameGPU; }; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 161283db2a2bc..3e50aedab5323 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -13,11 +13,13 @@ #ifndef ITSTRACKINGGPU_TRACKINGKERNELS_H_ #define ITSTRACKINGGPU_TRACKINGKERNELS_H_ +#include #include #include "ITStracking/BoundedAllocator.h" #include "ITStracking/ROFLookupTables.h" #include "ITStracking/TrackingTopology.h" +#include "ITStracking/TrackExtensionCandidate.h" #include "ITStrackingGPU/Utils.h" #include "DetectorsBase/Propagator.h" @@ -35,6 +37,58 @@ class Cluster; class TrackITSExt; class ExternalAllocator; +inline constexpr int kTrackExtensionLaunchBlocks = 60; +inline constexpr int kTrackExtensionLaunchThreadsPerBlock = 256; +inline constexpr int kTrackExtensionLaunchThreads = kTrackExtensionLaunchBlocks * kTrackExtensionLaunchThreadsPerBlock; + +template +void computeTrackExtensionCandidatesHandler(const TrackITSExt* tracks, + const IndexTableUtils* utils, + const typename ROFMaskTable::View& rofMask, + const typename ROFOverlapTable::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, + const TrackingFrameInfo** trackingFrameInfo, + TrackExtensionCandidate* candidates, + int* candidateOffsets, + TrackExtensionHypothesis* activeHypotheses, + TrackExtensionHypothesis* nextHypotheses, + const std::array layerRadii, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const float nSigmaCutPhi, + const float nSigmaCutZ, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + gpu::Stream& stream); + +template +void computeTrackExtensionResultsHandler(const TrackITSExt* tracks, + const TrackExtensionCandidate* candidates, + const int* candidateOffsets, + TrackExtensionResult* results, + const TrackingFrameInfo** trackingFrameInfo, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + gpu::Stream& stream); + template void countTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index 5fff30f5162b1..af6a86665de96 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -12,6 +12,7 @@ #include +#include #include #include @@ -581,6 +582,72 @@ void TimeFrameGPU::createTrackITSExtDevice(const size_t nSeeds) GPUChkErrS(cudaMemset(mTrackITSExtDevice, 0, mNTracks * sizeof(o2::its::TrackITSExt))); } +template +void TimeFrameGPU::loadTrackExtensionStartTracksDevice() +{ + GPUTimer timer("loading track extension start tracks"); + GPULog("gpu-transfer: loading {} track extension start tracks, for {:.2f} MB.", this->mTracks.size(), this->mTracks.size() * sizeof(o2::its::TrackITSExt) / constants::MB); + mTrackExtensionStartTracksDevice = nullptr; + mTrackExtensionStartTracks = bounded_vector(this->mTracks.begin(), this->mTracks.end(), this->getMemoryPool().get()); + if (this->mTracks.empty()) { + return; + } + allocMem(reinterpret_cast(&mTrackExtensionStartTracksDevice), mTrackExtensionStartTracks.size() * sizeof(o2::its::TrackITSExt), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); + GPUChkErrS(cudaMemcpy(mTrackExtensionStartTracksDevice, mTrackExtensionStartTracks.data(), mTrackExtensionStartTracks.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyHostToDevice)); +} + +template +void TimeFrameGPU::createTrackExtensionCandidatesDevice(const size_t nTracks) +{ + GPUTimer timer("reserving track extension candidates"); + const size_t nCandidates = nTracks * MaxTrackExtensionCandidatesPerTrack; + GPULog("gpu-allocation: reserving {} track extension candidates, for {:.2f} MB.", nCandidates, nCandidates * sizeof(o2::its::TrackExtensionCandidate) / constants::MB); + mTrackExtensionCandidatesDevice = nullptr; + mTrackExtensionCandidateOffsetsDevice = nullptr; + if (nCandidates == 0) { + return; + } + allocMem(reinterpret_cast(&mTrackExtensionCandidatesDevice), nCandidates * sizeof(o2::its::TrackExtensionCandidate), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); + allocMem(reinterpret_cast(&mTrackExtensionCandidateOffsetsDevice), (nTracks + 1) * sizeof(int), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); +} + +template +void TimeFrameGPU::createTrackExtensionScratchDevice(const int nThreads, const int beamWidth) +{ + GPUTimer timer("reserving track extension scratch"); + const size_t nHypotheses = static_cast(std::max(1, nThreads)) * std::max(1, beamWidth); + GPULog("gpu-allocation: reserving {} track extension hypotheses per scratch buffer, for {:.2f} MB each.", nHypotheses, nHypotheses * sizeof(o2::its::TrackExtensionHypothesis) / constants::MB); + mActiveTrackExtensionHypothesesDevice = nullptr; + mNextTrackExtensionHypothesesDevice = nullptr; + if (nHypotheses == 0) { + return; + } + allocMem(reinterpret_cast(&mActiveTrackExtensionHypothesesDevice), nHypotheses * sizeof(o2::its::TrackExtensionHypothesis), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); + allocMem(reinterpret_cast(&mNextTrackExtensionHypothesesDevice), nHypotheses * sizeof(o2::its::TrackExtensionHypothesis), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); +} + +template +void TimeFrameGPU::createTrackExtensionResultsDevice(const size_t nTracks) +{ + GPUTimer timer("reserving fitted track extension results"); + mNTrackExtensionResults = 0; + if (nTracks == 0 || mTrackExtensionCandidateOffsetsDevice == nullptr) { + mTrackExtensionResults = bounded_vector>(0, {}, this->getMemoryPool().get()); + mTrackExtensionResultsDevice = nullptr; + return; + } + int nResults{0}; + GPUChkErrS(cudaMemcpy(&nResults, mTrackExtensionCandidateOffsetsDevice + nTracks, sizeof(int), cudaMemcpyDeviceToHost)); + mNTrackExtensionResults = nResults; + GPULog("gpu-allocation: reserving {} fitted track extension results, for {:.2f} MB.", mNTrackExtensionResults, mNTrackExtensionResults * sizeof(o2::its::TrackExtensionResult) / constants::MB); + mTrackExtensionResults = bounded_vector>(mNTrackExtensionResults, {}, this->getMemoryPool().get()); + mTrackExtensionResultsDevice = nullptr; + if (mTrackExtensionResults.empty()) { + return; + } + allocMem(reinterpret_cast(&mTrackExtensionResultsDevice), mNTrackExtensionResults * sizeof(o2::its::TrackExtensionResult), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); +} + template void TimeFrameGPU::downloadCellsDevice() { @@ -627,6 +694,17 @@ void TimeFrameGPU::downloadTrackITSExtDevice() GPUChkErrS(cudaMemcpy(mTrackITSExt.data(), mTrackITSExtDevice, mTrackITSExt.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyDeviceToHost)); } +template +void TimeFrameGPU::downloadTrackExtensionResultsDevice() +{ + GPUTimer timer("downloading fitted track extension results"); + GPULog("gpu-transfer: downloading {} fitted track extension results, for {:.2f} MB.", mTrackExtensionResults.size(), mTrackExtensionResults.size() * sizeof(o2::its::TrackExtensionResult) / constants::MB); + if (mTrackExtensionResults.empty()) { + return; + } + GPUChkErrS(cudaMemcpy(mTrackExtensionResults.data(), mTrackExtensionResultsDevice, mTrackExtensionResults.size() * sizeof(o2::its::TrackExtensionResult), cudaMemcpyDeviceToHost)); +} + template void TimeFrameGPU::unregisterHostMemory(const int maxLayers) { diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 141d558712e6d..b88b63d04e053 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -12,6 +12,9 @@ #include +#include +#include + #include "ITStrackingGPU/TrackerTraitsGPU.h" #include "ITStrackingGPU/TrackingKernels.h" #include "ITStracking/Configuration.h" @@ -390,10 +393,127 @@ void TrackerTraitsGPU::findRoads(const int iteration) mTimeFrameGPU->loadUsedClustersDevice(); } this->markTracks(iteration); - // wipe the artefact memory - mTimeFrameGPU->popMemoryStack(iteration); + if (!hasTrackFollower(iteration)) { + // wipe the artefact memory + mTimeFrameGPU->popMemoryStack(iteration); + } }; +template +void TrackerTraitsGPU::extendTracks(const int iteration) +{ + TrackerTraits::extendTracks(iteration); + mTimeFrameGPU->loadUsedClustersDevice(); + // wipe the artefact memory kept alive for the track follower + mTimeFrameGPU->popMemoryStack(iteration); +} + +template +bool TrackerTraitsGPU::hasTrackFollower(const int iteration) const +{ + return this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || + this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; +} + +template +void TrackerTraitsGPU::buildTrackExtensionCandidates(const int iteration, typename TrackerTraits::TrackExtensionCandidates& candidatesPerTrack) +{ + const auto nTracks = this->mTimeFrame->getTracks().size(); + const int beamWidth = std::max(1, this->mTrkParams[iteration].TrackFollowerBeamWidth); + mTimeFrameGPU->syncStreams(); + mTimeFrameGPU->loadTrackExtensionStartTracksDevice(); + mTimeFrameGPU->createTrackExtensionCandidatesDevice(nTracks); + mTimeFrameGPU->createTrackExtensionScratchDevice(kTrackExtensionLaunchThreads, beamWidth); + std::array layerRadii{}; + std::array layerxX0{}; + for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) { + layerRadii[iLayer] = this->mTrkParams[iteration].LayerRadii[iLayer]; + layerxX0[iLayer] = this->mTrkParams[iteration].LayerxX0[iLayer]; + } + computeTrackExtensionCandidatesHandler(mTimeFrameGPU->getDeviceTrackExtensionStartTracks(), + mTimeFrameGPU->getDeviceIndexTableUtils(), + mTimeFrameGPU->getDeviceROFMaskTableView(), + mTimeFrameGPU->getDeviceROFOverlapTableView(), + mTimeFrameGPU->getDeviceArrayClusters(), + (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(), + mTimeFrameGPU->getDeviceArrayClustersIndexTables(), + mTimeFrameGPU->getDeviceROFrameClusters(), + mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), + mTimeFrameGPU->getDeviceTrackExtensionCandidates(), + mTimeFrameGPU->getDeviceTrackExtensionCandidateOffsets(), + mTimeFrameGPU->getDeviceActiveTrackExtensionHypotheses(), + mTimeFrameGPU->getDeviceNextTrackExtensionHypotheses(), + layerRadii, + layerxX0, + static_cast(nTracks), + this->mTrkParams[iteration].NLayers, + this->mTrkParams[iteration].PhiBins, + beamWidth, + this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop], + this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot], + this->mBz, + this->mTrkParams[iteration].MaxChi2ClusterAttachment, + this->mTrkParams[iteration].MaxChi2NDF, + this->mTrkParams[iteration].TrackFollowerNSigmaCutPhi, + this->mTrkParams[iteration].TrackFollowerNSigmaCutZ, + mTimeFrameGPU->getDevicePropagator(), + this->mTrkParams[iteration].CorrType, + mTimeFrameGPU->getStream(0)); + mTimeFrameGPU->createTrackExtensionResultsDevice(nTracks); + computeTrackExtensionResultsHandler(mTimeFrameGPU->getDeviceTrackExtensionStartTracks(), + mTimeFrameGPU->getDeviceTrackExtensionCandidates(), + mTimeFrameGPU->getDeviceTrackExtensionCandidateOffsets(), + mTimeFrameGPU->getDeviceTrackExtensionResults(), + mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), + layerxX0, + static_cast(nTracks), + this->mTrkParams[iteration].NLayers, + this->mBz, + this->mTrkParams[iteration].MaxChi2ClusterAttachment, + this->mTrkParams[iteration].MaxChi2NDF, + mTimeFrameGPU->getDevicePropagator(), + this->mTrkParams[iteration].CorrType, + this->mTrkParams[iteration].ShiftRefToCluster, + mTimeFrameGPU->getStream(0)); + mTimeFrameGPU->downloadTrackExtensionResultsDevice(); + + const auto& results = mTimeFrameGPU->getTrackExtensionResults(); + for (int iResult{0}; iResult < static_cast(results.size()); ++iResult) { + const auto& result = results[iResult]; + if (!result.isValid()) { + continue; + } + auto candidate = result.candidate; + candidate.resultIndex = iResult; + if (candidatesPerTrack.add(candidate.trackIndex, candidate) < 0) { + continue; + } + } +} + +template +bool TrackerTraitsGPU::materializeTrackExtensionCandidate(TrackITSExt& track, const typename TrackerTraits::TrackExtensionCandidateN& candidate, const int iteration) +{ + const auto& results = mTimeFrameGPU->getTrackExtensionResults(); + if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast(results.size())) { + return TrackerTraits::materializeTrackExtensionCandidate(track, candidate, iteration); + } + const auto& result = results[candidate.resultIndex]; + if (!result.isValid() || result.candidate.trackIndex != candidate.trackIndex) { + return false; + } + track = result.track; + this->updateExtendedTrackTimeStamp(track, iteration); + uint32_t diff{0}; + for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) { + if (candidate.addedClusters[iLayer] != constants::UnusedIndex) { + diff |= (0x1u << iLayer); + } + } + applyExtendedClustersPattern(track, diff); + return true; +} + template int TrackerTraitsGPU::getTFNumberOfClusters() const { diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 571afe08fc209..d42b373ca3e64 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -20,6 +20,7 @@ #include #include #include +#include #include #include @@ -31,6 +32,7 @@ #include "ITStracking/Tracklet.h" #include "ITStracking/Cluster.h" #include "ITStracking/Cell.h" +#include "ITStracking/TrackFollower.h" #include "ITStracking/TrackHelpers.h" #include "DataFormatsITS/TrackITS.h" #include "ITStrackingGPU/TrackingKernels.h" @@ -108,6 +110,229 @@ struct compare_track_chi2 { } }; +template +GPUdi() void writeTrackExtensionCandidate(const int trackIndex, + const TrackITSExt& original, + const TrackITSExt& updated, + TrackExtensionCandidate* candidates, + int& slot) +{ + if (slot >= MaxTrackExtensionCandidatesPerTrack) { + return; + } + auto& candidate = candidates[getFlatTrackExtensionCandidateIndex(trackIndex, slot)]; + candidate.reset(); + candidate.trackIndex = trackIndex; + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + if (original.getClusterIndex(iLayer) == constants::UnusedIndex && updated.getClusterIndex(iLayer) != constants::UnusedIndex) { + candidate.addedClusters[iLayer] = updated.getClusterIndex(iLayer); + ++candidate.nAddedClusters; + } + } + if (!candidate.nAddedClusters) { + candidate.reset(); + return; + } + candidate.chi2 = updated.getChi2(); + ++slot; +} + +template +GPUg() void __launch_bounds__(256, 1) computeTrackExtensionCandidatesKernel(const TrackITSExt* tracks, + const IndexTableUtils* utils, + const typename ROFMaskTable::View rofMask, + const typename ROFOverlapTable::View rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, + const TrackingFrameInfo** trackingFrameInfo, + TrackExtensionCandidate* candidates, + int* candidateOffsets, + TrackExtensionHypothesis* activeHypothesesScratch, + TrackExtensionHypothesis* nextHypothesesScratch, + const std::array layerRadii, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const float nSigmaCutPhi, + const float nSigmaCutZ, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType) +{ + if (blockIdx.x == 0 && threadIdx.x == 0) { + candidateOffsets[nTracks] = 0; + } + const int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x; + auto* const threadActiveHypotheses = activeHypothesesScratch + (globalThreadId * beamWidth); + auto* const threadNextHypotheses = nextHypothesesScratch + (globalThreadId * beamWidth); + for (int iTrack = globalThreadId; iTrack < nTracks; iTrack += blockDim.x * gridDim.x) { + for (int iCandidate{0}; iCandidate < MaxTrackExtensionCandidatesPerTrack; ++iCandidate) { + candidates[getFlatTrackExtensionCandidateIndex(iTrack, iCandidate)].reset(); + } + const auto& track = tracks[iTrack]; + auto* activeHypotheses = threadActiveHypotheses; + auto* nextHypotheses = threadNextHypotheses; + int slot{0}; + if (extendTop && getTrackExtensionLastClusterLayer(track) != nLayers - 1) { + TrackITSExt topCandidate; + if (followTrackExtensionDirection(track, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, true, propagator, matCorrType, activeHypotheses, nextHypotheses, topCandidate)) { + writeTrackExtensionCandidate(iTrack, track, topCandidate, candidates, slot); + if (extendBot && getTrackExtensionFirstClusterLayer(topCandidate) != 0) { + TrackITSExt topBottomCandidate; + if (followTrackExtensionDirection(topCandidate, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, false, propagator, matCorrType, activeHypotheses, nextHypotheses, topBottomCandidate)) { + writeTrackExtensionCandidate(iTrack, track, topBottomCandidate, candidates, slot); + } + } + } + } + if (extendBot && getTrackExtensionFirstClusterLayer(track) != 0) { + TrackITSExt bottomCandidate; + if (followTrackExtensionDirection(track, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, false, propagator, matCorrType, activeHypotheses, nextHypotheses, bottomCandidate)) { + writeTrackExtensionCandidate(iTrack, track, bottomCandidate, candidates, slot); + if (extendTop && getTrackExtensionLastClusterLayer(bottomCandidate) != nLayers - 1) { + TrackITSExt bottomTopCandidate; + if (followTrackExtensionDirection(bottomCandidate, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, true, propagator, matCorrType, activeHypotheses, nextHypotheses, bottomTopCandidate)) { + writeTrackExtensionCandidate(iTrack, track, bottomTopCandidate, candidates, slot); + } + } + } + } + candidateOffsets[iTrack] = slot; + } +} + +template +GPUdi() bool fitTrackExtensionResult(const TrackITSExt& startTrack, + const TrackExtensionCandidate& candidate, + const TrackingFrameInfo* const* trackingFrameInfo, + const float* layerxX0, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + TrackITSExt& track) +{ + track = startTrack; + for (int iLayer{0}; iLayer < nLayers; ++iLayer) { + if (candidate.addedClusters[iLayer] != constants::UnusedIndex) { + track.setExternalClusterIndex(iLayer, candidate.addedClusters[iLayer], true); + } + } + + o2::track::TrackPar linRef{track}; + o2::its::track::resetTrackCovariance(track); + track.setChi2(0); + bool fitSuccess = o2::its::track::fitTrack(track, + 0, + nLayers, + 1, + maxChi2ClusterAttachment, + maxChi2NDF, + o2::constants::math::VeryBig, + 0, + bz, + trackingFrameInfo, + layerxX0, + propagator, + matCorrType, + &linRef, + shiftRefToCluster); + if (!fitSuccess) { + return false; + } + + track.getParamOut() = track.getParamIn(); + linRef = track.getParamOut(); + o2::its::track::resetTrackCovariance(track); + track.setChi2(0); + fitSuccess = o2::its::track::fitTrack(track, + nLayers - 1, + -1, + -1, + maxChi2ClusterAttachment, + maxChi2NDF, + 50.f, + 0, + bz, + trackingFrameInfo, + layerxX0, + propagator, + matCorrType, + &linRef, + shiftRefToCluster); + if (!fitSuccess) { + return false; + } + + uint32_t diff{0}; + for (int iLayer{0}; iLayer < nLayers; ++iLayer) { + if (candidate.addedClusters[iLayer] != constants::UnusedIndex) { + diff |= (0x1u << iLayer); + } + } + applyExtendedClustersPattern(track, diff); + return true; +} + +template +GPUg() void __launch_bounds__(256, 1) computeTrackExtensionResultsKernel(const TrackITSExt* tracks, + const TrackExtensionCandidate* candidates, + const int* candidateOffsets, + TrackExtensionResult* results, + const TrackingFrameInfo** trackingFrameInfo, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster) +{ + for (int iTrack = blockIdx.x * blockDim.x + threadIdx.x; iTrack < nTracks; iTrack += blockDim.x * gridDim.x) { + const int firstResult = candidateOffsets[iTrack]; + const int nResults = candidateOffsets[iTrack + 1] - firstResult; + const auto& startTrack = tracks[iTrack]; + for (int iCandidate{0}; iCandidate < nResults; ++iCandidate) { + const auto& candidate = candidates[getFlatTrackExtensionCandidateIndex(iTrack, iCandidate)]; + auto& result = results[firstResult + iCandidate]; + result.reset(); + if (!candidate.isValidForTrack(iTrack)) { + continue; + } + result.candidate = candidate; + if (!fitTrackExtensionResult(startTrack, + candidate, + trackingFrameInfo, + layerxX0.data(), + nLayers, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + propagator, + matCorrType, + shiftRefToCluster, + result.track)) { + result.reset(); + continue; + } + result.candidate.chi2 = result.track.getChi2(); + } + } +} + template GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( TrackSeed* trackSeeds, @@ -584,6 +809,114 @@ GPUg() void __launch_bounds__(256, 1) processNeighboursKernel( } // namespace gpu +template +void computeTrackExtensionCandidatesHandler(const TrackITSExt* tracks, + const IndexTableUtils* utils, + const typename ROFMaskTable::View& rofMask, + const typename ROFOverlapTable::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, + const TrackingFrameInfo** trackingFrameInfo, + TrackExtensionCandidate* candidates, + int* candidateOffsets, + TrackExtensionHypothesis* activeHypotheses, + TrackExtensionHypothesis* nextHypotheses, + const std::array layerRadii, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const float nSigmaCutPhi, + const float nSigmaCutZ, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + gpu::Stream& stream) +{ + if (nTracks <= 0 || candidates == nullptr || candidateOffsets == nullptr || activeHypotheses == nullptr || nextHypotheses == nullptr) { + return; + } + gpu::computeTrackExtensionCandidatesKernel<<>>( + tracks, + utils, + rofMask, + rofOverlaps, + clusters, + usedClusters, + clustersIndexTables, + ROFClusters, + trackingFrameInfo, + candidates, + candidateOffsets, + activeHypotheses, + nextHypotheses, + layerRadii, + layerxX0, + nTracks, + nLayers, + phiBins, + beamWidth, + extendTop, + extendBot, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + nSigmaCutPhi, + nSigmaCutZ, + propagator, + matCorrType); + GPUChkErrS(cudaGetLastError()); + GPUChkErrS(cudaStreamSynchronize(stream.get())); + thrust::device_ptr offsets(candidateOffsets); + thrust::exclusive_scan(offsets, offsets + nTracks + 1, offsets); +} + +template +void computeTrackExtensionResultsHandler(const TrackITSExt* tracks, + const TrackExtensionCandidate* candidates, + const int* candidateOffsets, + TrackExtensionResult* results, + const TrackingFrameInfo** trackingFrameInfo, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + gpu::Stream& stream) +{ + if (nTracks <= 0 || tracks == nullptr || candidates == nullptr || candidateOffsets == nullptr || results == nullptr) { + return; + } + gpu::computeTrackExtensionResultsKernel<<>>( + tracks, + candidates, + candidateOffsets, + results, + trackingFrameInfo, + layerxX0, + nTracks, + nLayers, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + propagator, + matCorrType, + shiftRefToCluster); + GPUChkErrS(cudaGetLastError()); + GPUChkErrS(cudaStreamSynchronize(stream.get())); +} + template void countTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, @@ -1131,6 +1464,52 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, } /// Explicit instantiation of ITS2 handlers +template void computeTrackExtensionCandidatesHandler<7>(const TrackITSExt* tracks, + const IndexTableUtils<7>* utils, + const ROFMaskTable<7>::View& rofMask, + const ROFOverlapTable<7>::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, + const TrackingFrameInfo** trackingFrameInfo, + TrackExtensionCandidate<7>* candidates, + int* candidateOffsets, + TrackExtensionHypothesis<7>* activeHypotheses, + TrackExtensionHypothesis<7>* nextHypotheses, + const std::array layerRadii, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const float nSigmaCutPhi, + const float nSigmaCutZ, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + gpu::Stream& stream); + +template void computeTrackExtensionResultsHandler<7>(const TrackITSExt* tracks, + const TrackExtensionCandidate<7>* candidates, + const int* candidateOffsets, + TrackExtensionResult<7>* results, + const TrackingFrameInfo** trackingFrameInfo, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + gpu::Stream& stream); + template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const ROFMaskTable<7>::View& rofMask, const int transitionId, @@ -1317,6 +1696,52 @@ template void computeTrackSeedHandler(TrackSeed<7>* trackSeeds, /// Explicit instantiation of ALICE3 handlers #ifdef ENABLE_UPGRADES +template void computeTrackExtensionCandidatesHandler<11>(const TrackITSExt* tracks, + const IndexTableUtils<11>* utils, + const ROFMaskTable<11>::View& rofMask, + const ROFOverlapTable<11>::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, + const TrackingFrameInfo** trackingFrameInfo, + TrackExtensionCandidate<11>* candidates, + int* candidateOffsets, + TrackExtensionHypothesis<11>* activeHypotheses, + TrackExtensionHypothesis<11>* nextHypotheses, + const std::array layerRadii, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const float nSigmaCutPhi, + const float nSigmaCutZ, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + gpu::Stream& stream); + +template void computeTrackExtensionResultsHandler<11>(const TrackITSExt* tracks, + const TrackExtensionCandidate<11>* candidates, + const int* candidateOffsets, + TrackExtensionResult<11>* results, + const TrackingFrameInfo** trackingFrameInfo, + const std::array layerxX0, + const int nTracks, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + gpu::Stream& stream); + template void countTrackletsInROFsHandler<11>(const IndexTableUtils<11>* utils, const ROFMaskTable<11>::View& rofMask, const int transitionId, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 275752854665b..5c1dcf5216f51 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -32,7 +32,7 @@ namespace o2::its { // Steering of dedicated steps in an iteration -enum class IterationStep : uint8_t { +enum class IterationStep : uint16_t { FirstPass = 0, RebuildClusterLUT, UseUPCMask, @@ -40,6 +40,8 @@ enum class IterationStep : uint8_t { ResetVertices, SkipROFsAboveThreshold, MarkVerticesAsUPC, + TrackFollowerTop, + TrackFollowerBot, }; using IterationSteps = o2::utils::EnumFlags; @@ -94,6 +96,9 @@ struct TrackingParameters { bool DoUPCIteration = false; bool FataliseUponFailure = true; bool CreateArtefactLabels{false}; + float TrackFollowerNSigmaCutZ = 1.f; + float TrackFollowerNSigmaCutPhi = 1.f; + int TrackFollowerBeamWidth = 1; bool PrintMemory = false; // print allocator usage in epilog report size_t MaxMemory = std::numeric_limits::max(); bool DropTFUponFailure = false; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h index ce20169e36c64..a8e2c37e261fb 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h @@ -73,7 +73,7 @@ struct LayerTiming { } // return which ROF this BC belongs to - GPUhi() BCType getROF(BCType bc) const noexcept + GPUhdi() BCType getROF(BCType bc) const noexcept { const BCType offset = mROFDelay + mROFBias; if (bc <= offset) { @@ -83,7 +83,7 @@ struct LayerTiming { } // return which ROF this timestamp belongs by its lower edge - GPUhi() BCType getROF(TimeStamp ts) const noexcept + GPUhdi() BCType getROF(TimeStamp ts) const noexcept { const BCType offset = mROFDelay + mROFBias; const BCType bc = (ts.getTimeStamp() < ts.getTimeStampError()) ? BCType(0) : static_cast(o2::gpu::CAMath::Floor(ts.getTimeStamp() - ts.getTimeStampError())); @@ -93,6 +93,45 @@ struct LayerTiming { return (bc - offset) / mROFLength; } + // return which ROF this floating point (number of BCs) time belongs + GPUhdi() BCType getROF(float time) const noexcept + { + const float offset = static_cast(mROFDelay + mROFBias); + if (time <= offset) { + return 0; + } + return static_cast((time - offset) / mROFLength); + } + + GPUhdi() bool intersectROF(BCType rof, float lower, float upper) const noexcept + { + const auto rofTS = getROFTimeBounds(rof, true); + return static_cast(rofTS.upper()) > lower && upper > static_cast(rofTS.lower()); + } + + // return clamped ROF range with strictly positive overlap with timestamp interval + GPUhdi() int2 getROFRange(TimeStamp ts) const noexcept + { + if (mNROFsTF == 0) { + return {1, 0}; + } + + const float lower = ts.getTimeStamp() - ts.getTimeStampError(); + const float upper = ts.getTimeStamp() + ts.getTimeStampError(); + const int maxROF = static_cast(mNROFsTF) - 1; + int2 range{ + o2::gpu::CAMath::Clamp(static_cast(getROF(lower - mROFAddTimeErr)), 0, maxROF), + o2::gpu::CAMath::Clamp(static_cast(getROF(upper + mROFAddTimeErr)), 0, maxROF)}; + + if (range.x <= range.y && !intersectROF(static_cast(range.x), lower, upper)) { + ++range.x; + } + if (range.y >= range.x && !intersectROF(static_cast(range.y), lower, upper)) { + --range.y; + } + return range; + } + #ifndef GPUCA_GPUCODE GPUh() std::string asString() const { diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index 3fef2dc640cbc..7d908d6265660 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -279,6 +279,7 @@ struct TimeFrame { std::vector> mTracklets; std::vector> mCells; bounded_vector mTracks; + std::vector mFittedExtensionTracks; bounded_vector mTracksLabel; std::vector> mCellsNeighbours; std::vector> mCellsNeighboursTopology; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h new file mode 100644 index 0000000000000..be165b54df8c6 --- /dev/null +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h @@ -0,0 +1,114 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#ifndef TRACKINGITSU_INCLUDE_TRACKEXTENSIONCANDIDATE_H_ +#define TRACKINGITSU_INCLUDE_TRACKEXTENSIONCANDIDATE_H_ + +#include +#include + +#include "GPUCommonDef.h" +#include "DataFormatsITS/TrackITS.h" +#include "DataFormatsITS/TimeEstBC.h" +#include "ITStracking/Constants.h" +#include "ReconstructionDataFormats/Track.h" + +namespace o2::its +{ + +inline constexpr unsigned int kExtendedPatternShift = 24; +inline constexpr int kMaxLayersInTrackPattern = 8; + +template +GPUhdi() constexpr uint32_t makeAddedClustersPatternMask() +{ + return (NLayers >= 32) ? 0xffffffffu : ((1u << NLayers) - 1u); +} + +template +GPUhdi() void applyExtendedClustersPattern(TrackITSExt& track, uint32_t diff) +{ + if constexpr (NLayers <= kMaxLayersInTrackPattern) { + track.setPattern(track.getPattern() | (diff << kExtendedPatternShift)); + } else { + (void)track; + (void)diff; + } +} + +template +struct TrackExtensionHypothesis { + o2::track::TrackParCov param; + std::array clusters{}; + TimeStamp time; + float chi2{0.f}; + int nClusters{0}; + int edgeLayer{constants::UnusedIndex}; +}; + +template +struct TrackExtensionCandidate { + static constexpr float InvalidChi2 = 1.e20f; + + GPUhdi() TrackExtensionCandidate() { reset(); } + + GPUhdi() void reset() + { + trackIndex = -1; + nAddedClusters = 0; + resultIndex = -1; + chi2 = InvalidChi2; + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + addedClusters[iLayer] = constants::UnusedIndex; + } + } + + GPUhdi() bool isValidForTrack(int index) const + { + return trackIndex == index && nAddedClusters > 0; + } + + int trackIndex{-1}; + std::array addedClusters; + int nAddedClusters{0}; + int resultIndex{-1}; + float chi2{InvalidChi2}; +}; + +template +GPUhdi() bool isBetterTrackExtensionCandidate(const TrackExtensionCandidate& a, const TrackExtensionCandidate& b) +{ + return (a.nAddedClusters > b.nAddedClusters) || (a.nAddedClusters == b.nAddedClusters && a.chi2 < b.chi2); +} + +template +struct TrackExtensionResult { + GPUhdi() void reset() + { + candidate.reset(); + } + + GPUhdi() bool isValid() const { return candidate.trackIndex >= 0 && candidate.nAddedClusters > 0; } + + TrackExtensionCandidate candidate; + TrackITSExt track; +}; + +inline constexpr int MaxTrackExtensionCandidatesPerTrack = 4; + +inline constexpr size_t getFlatTrackExtensionCandidateIndex(size_t trackIndex, size_t candidateIndex) +{ + return trackIndex * MaxTrackExtensionCandidatesPerTrack + candidateIndex; +} + +} // namespace o2::its + +#endif /* TRACKINGITSU_INCLUDE_TRACKEXTENSIONCANDIDATE_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h new file mode 100644 index 0000000000000..8cd20262edf14 --- /dev/null +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h @@ -0,0 +1,301 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file TrackFollower.h +/// \brief Beam search used by CPU and GPU track extension. + +#ifndef TRACKINGITSU_INCLUDE_TRACKFOLLOWER_H_ +#define TRACKINGITSU_INCLUDE_TRACKFOLLOWER_H_ + +#include "GPUCommonDef.h" +#include "GPUCommonMath.h" +#include "CommonConstants/MathConstants.h" +#include "DetectorsBase/Propagator.h" + +#include "ITStracking/Cluster.h" +#include "ITStracking/Constants.h" +#include "ITStracking/IndexTableUtils.h" +#include "ITStracking/MathUtils.h" +#include "ITStracking/ROFLookupTables.h" +#include "ITStracking/TrackExtensionCandidate.h" + +namespace o2::its +{ + +template +GPUhdi() bool isBetterTrackExtensionHypothesis(const TrackExtensionHypothesis& a, const TrackExtensionHypothesis& b) +{ + return (a.nClusters > b.nClusters) || (a.nClusters == b.nClusters && a.chi2 < b.chi2); +} + +template +GPUhdi() void addTrackExtensionHypothesisToBeam(const TrackExtensionHypothesis& hypo, + TrackExtensionHypothesis* beam, + int& nBeam, + const int beamWidth) +{ + if (nBeam < beamWidth) { + beam[nBeam++] = hypo; + return; + } + + int worst{0}; + for (int i{1}; i < nBeam; ++i) { + if (isBetterTrackExtensionHypothesis(beam[worst], beam[i])) { + worst = i; + } + } + if (isBetterTrackExtensionHypothesis(hypo, beam[worst])) { + beam[worst] = hypo; + } +} + +template +GPUhdi() int4 getTrackExtensionBinsAt(const IndexTableUtils& utils, + const int layer, + const float phi, + const float deltaPhi, + const float z, + const float deltaZ) +{ + const float zRangeMin = z - deltaZ; + const float zRangeMax = z + deltaZ; + if (zRangeMax < -utils.getLayerZ(layer) || zRangeMin > utils.getLayerZ(layer) || zRangeMin > zRangeMax) { + return {-1, -1, -1, -1}; + } + const float phiRangeMin = (deltaPhi > o2::constants::math::PI) ? 0.f : phi - deltaPhi; + const float phiRangeMax = (deltaPhi > o2::constants::math::PI) ? o2::constants::math::TwoPI : phi + deltaPhi; + return {o2::gpu::CAMath::Max(0, utils.getZBinIndex(layer, zRangeMin)), + utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMin)), + o2::gpu::CAMath::Min(utils.getNzBins() - 1, utils.getZBinIndex(layer, zRangeMax)), + utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))}; +} + +template +GPUhdi() int getTrackExtensionFirstClusterLayer(const TrackITSExt& track) +{ + const uint32_t pattern = track.getPattern(); + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + if (pattern & (0x1u << iLayer)) { + return iLayer; + } + } + return constants::UnusedIndex; +} + +template +GPUhdi() int getTrackExtensionLastClusterLayer(const TrackITSExt& track) +{ + const uint32_t pattern = track.getPattern(); + for (int iLayer{NLayers}; iLayer-- > 0;) { + if (pattern & (0x1u << iLayer)) { + return iLayer; + } + } + return constants::UnusedIndex; +} + +template +GPUhdi() void initialiseTrackExtensionHypothesis(const TrackITSExt& track, + const bool outward, + TrackExtensionHypothesis& hypo) +{ + hypo.param = outward ? track.getParamOut() : track.getParamIn(); + hypo.time = track.getTimeStamp(); + hypo.chi2 = track.getChi2(); + hypo.nClusters = track.getNClusters(); + hypo.edgeLayer = outward ? getTrackExtensionLastClusterLayer(track) : getTrackExtensionFirstClusterLayer(track); + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + hypo.clusters[iLayer] = track.getClusterIndex(iLayer); + } +} + +template +GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, + const IndexTableUtils& utils, + const typename ROFMaskTable::View& rofMask, + const typename ROFOverlapTable::View& rofOverlaps, + const Cluster* const* clusters, + const unsigned char* const* usedClusters, + const int* const* clustersIndexTables, + const int* const* ROFClusters, + const TrackingFrameInfo* const* trackingFrameInfo, + const float* layerRadii, + const float* layerxX0, + const int nLayers, + const int phiBins, + const int beamWidthConfig, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const float nSigmaCutPhi, + const float nSigmaCutZ, + const bool outward, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + TrackExtensionHypothesis* activeHypotheses, + TrackExtensionHypothesis* nextHypotheses, + TrackITSExt& updatedTrack) +{ + const int step = outward ? 1 : -1; + const int end = outward ? nLayers - 1 : 0; + const int beamWidth = o2::gpu::CAMath::Max(beamWidthConfig, 1); + int nActive{1}; + int nNext{0}; + initialiseTrackExtensionHypothesis(track, outward, activeHypotheses[0]); + + const int tableSize = utils.getNphiBins() * utils.getNzBins() + 1; + for (int iLayer = activeHypotheses[0].edgeLayer + step; nActive > 0; iLayer += step) { + if ((step > 0 && iLayer > end) || (step < 0 && iLayer < end)) { + break; + } + nNext = 0; + for (int iHypo{0}; iHypo < nActive; ++iHypo) { + auto hypo = activeHypotheses[iHypo]; + const float r = layerRadii[iLayer]; + float x{-999.f}; + if (!hypo.param.getXatLabR(r, x, bz, o2::track::DirAuto) || x <= 0.f) { + continue; + } + + if (!propagator->propagateToX(hypo.param, x, bz, o2::base::PropagatorF::MAX_SIN_PHI, + o2::base::PropagatorF::MAX_STEP, matCorrType)) { + continue; + } + if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE && + !hypo.param.correctForMaterial(layerxX0[iLayer], layerxX0[iLayer] * constants::Radl * constants::Rho, true)) { + continue; + } + + const float ePhi{o2::gpu::CAMath::Sqrt(hypo.param.getSigmaSnp2() / hypo.param.getCsp2())}; + const float eZ{o2::gpu::CAMath::Sqrt(hypo.param.getSigmaZ2())}; + const int4 selectedBins = getTrackExtensionBinsAt(utils, + iLayer, + hypo.param.getPhi(), + nSigmaCutPhi * ePhi, + hypo.param.getZ(), + nSigmaCutZ * eZ); + if (selectedBins.x < 0) { + continue; + } + + int phiBinsNum = selectedBins.w - selectedBins.y + 1; + if (phiBinsNum < 0) { + phiBinsNum += phiBins; + } + + const auto rofRange = rofOverlaps.getLayer(iLayer).getROFRange(hypo.time); + for (int rof = rofRange.x; rof <= rofRange.y; ++rof) { + if (!rofMask.isROFEnabled(iLayer, rof)) { + continue; + } + const int rofStart = ROFClusters[iLayer][rof]; + const int nLayerClusters = ROFClusters[iLayer][rof + 1] - rofStart; + if (nLayerClusters <= 0) { + continue; + } + const Cluster* layerClusters = clusters[iLayer] + rofStart; + const int* indexTable = clustersIndexTables[iLayer] + rof * tableSize; + const int zBinRange = selectedBins.z - selectedBins.x + 1; + for (int iPhiCount = 0; iPhiCount < phiBinsNum; ++iPhiCount) { + const int iPhiBin = (selectedBins.y + iPhiCount) % phiBins; + const int firstBinIndex = utils.getBinIndex(selectedBins.x, iPhiBin); + const int maxBinIndex = firstBinIndex + zBinRange; + const int firstRowClusterIndex = indexTable[firstBinIndex]; + const int maxRowClusterIndex = indexTable[maxBinIndex]; + for (int iNextCluster{firstRowClusterIndex}; iNextCluster < maxRowClusterIndex; ++iNextCluster) { + if (iNextCluster >= nLayerClusters) { + break; + } + const Cluster& nextCluster = layerClusters[iNextCluster]; + if (usedClusters[iLayer][nextCluster.clusterId]) { + continue; + } + + const TrackingFrameInfo& trackingHit = trackingFrameInfo[iLayer][nextCluster.clusterId]; + auto updated = hypo; + if (!updated.param.rotate(trackingHit.alphaTrackingFrame) || + !propagator->propagateToX(updated.param, trackingHit.xTrackingFrame, bz, + o2::base::PropagatorF::MAX_SIN_PHI, + o2::base::PropagatorF::MAX_STEP, + matCorrType)) { + continue; + } + + const auto predChi2 = updated.param.getPredictedChi2Quiet(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame); + if (predChi2 < 0.f || predChi2 > maxChi2ClusterAttachment) { + continue; + } + if (!updated.param.o2::track::TrackParCov::update(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)) { + continue; + } + updated.chi2 += predChi2; + updated.clusters[iLayer] = nextCluster.clusterId; + ++updated.nClusters; + updated.edgeLayer = iLayer; + const auto rofTS = rofOverlaps.getLayer(iLayer).getROFTimeBounds(rof, true); + const auto& ts = updated.time; + const float lower = o2::gpu::CAMath::Max(ts.getTimeStamp() - ts.getTimeStampError(), static_cast(rofTS.lower())); + const float upper = o2::gpu::CAMath::Min(ts.getTimeStamp() + ts.getTimeStampError(), static_cast(rofTS.upper())); + updated.time.setTimeStamp(0.5f * (lower + upper)); + updated.time.setTimeStampError(0.5f * (upper - lower)); + addTrackExtensionHypothesisToBeam(updated, nextHypotheses, nNext, beamWidth); + } + } + } + addTrackExtensionHypothesisToBeam(hypo, nextHypotheses, nNext, beamWidth); + } + if (nNext == 0) { + break; + } + for (int iHypo{0}; iHypo < nNext; ++iHypo) { + activeHypotheses[iHypo] = nextHypotheses[iHypo]; + } + nActive = nNext; + } + + const TrackExtensionHypothesis* bestHypo{nullptr}; + for (int iHypo{0}; iHypo < nActive; ++iHypo) { + const auto& hypo = activeHypotheses[iHypo]; + if (hypo.nClusters == track.getNClusters()) { + continue; + } + const float maxChi2 = maxChi2NDF * static_cast(hypo.nClusters * 2 - 5); + if (hypo.chi2 >= maxChi2) { + continue; + } + if (!bestHypo || isBetterTrackExtensionHypothesis(hypo, *bestHypo)) { + bestHypo = &hypo; + } + } + if (!bestHypo) { + return false; + } + + updatedTrack = track; + if (outward) { + updatedTrack.getParamOut() = bestHypo->param; + } else { + updatedTrack.getParamIn() = bestHypo->param; + } + updatedTrack.getTimeStamp() = bestHypo->time; + updatedTrack.setChi2(bestHypo->chi2); + for (int iLayer{0}; iLayer < nLayers; ++iLayer) { + if (updatedTrack.getClusterIndex(iLayer) == constants::UnusedIndex && bestHypo->clusters[iLayer] != constants::UnusedIndex) { + updatedTrack.setExternalClusterIndex(iLayer, bestHypo->clusters[iLayer], true); + } + } + return true; +} + +} // namespace o2::its + +#endif // TRACKINGITSU_INCLUDE_TRACKFOLLOWER_H_ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h index 240b0eb1e2f63..daa185c945560 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h @@ -74,6 +74,7 @@ class Tracker void computeCells(int iteration) { mTraits->computeLayerCells(iteration); } void findCellsNeighbours(int iteration) { mTraits->findCellsNeighbours(iteration); } void findRoads(int iteration) { mTraits->findRoads(iteration); } + void extendTracks(int iteration) { mTraits->extendTracks(iteration); } void rectifyClusterIndices(); void sortTracks(); @@ -99,10 +100,11 @@ class Tracker Celling, Neighbouring, Roading, + Extending, NSteps, }; Steps mCurStep{TFInit}; - static constexpr std::array StateNames{"TimeFrame initialisation", "Tracklet finding", "Cell finding", "Neighbour finding", "Road finding"}; + static constexpr std::array StateNames{"TimeFrame initialisation", "Tracklet finding", "Cell finding", "Neighbour finding", "Road finding", "Track extending"}; std::vector> mTimingStats; void addTimingStatCurStep(int iteration, double timeMs); }; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index f536e86fe95d5..e870af0ad7baa 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -16,13 +16,19 @@ #ifndef TRACKINGITSU_INCLUDE_TRACKERTRAITS_H_ #define TRACKINGITSU_INCLUDE_TRACKERTRAITS_H_ +#include #include +#include #include "ITStracking/Configuration.h" +#include "ITStracking/Constants.h" #include "ITStracking/IndexTableUtils.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/Cell.h" #include "ITStracking/BoundedAllocator.h" +#include "DataFormatsITS/TimeEstBC.h" +#include "ReconstructionDataFormats/Track.h" +#include "ITStracking/TrackExtensionCandidate.h" // #define OPTIMISATION_OUTPUT @@ -51,6 +57,8 @@ class TrackerTraits virtual void computeLayerCells(const int iteration); virtual void findCellsNeighbours(const int iteration); virtual void findRoads(const int iteration); + virtual bool supportsExtendTracks() const noexcept { return true; } + virtual void extendTracks(const int iteration); template void processNeighbours(int iteration, int defaultCellTopologyId, int iLevel, const bounded_vector& currentCellSeed, const bounded_vector& currentCellId, const bounded_vector& currentCellTopologyId, bounded_vector& updatedCellSeed, bounded_vector& updatedCellId, bounded_vector& updatedCellTopologyId); @@ -85,6 +93,55 @@ class TrackerTraits std::shared_ptr mTaskArena; protected: + using TrackExtensionCandidateN = TrackExtensionCandidate; + struct TrackExtensionCandidates { + TrackExtensionCandidates() = default; + explicit TrackExtensionCandidates(size_t nTracks) + : candidates(nTracks * MaxTrackExtensionCandidatesPerTrack), counts(nTracks, 0) + { + } + + int add(int trackIndex, const TrackExtensionCandidateN& candidate) + { + auto& count = counts[trackIndex]; + if (count >= MaxTrackExtensionCandidatesPerTrack) { + return -1; + } + const int flatIndex = static_cast(getFlatTrackExtensionCandidateIndex(trackIndex, count)); + candidates[flatIndex] = candidate; + ++count; + return flatIndex; + } + + void pop_back(int trackIndex) + { + --counts[trackIndex]; + } + + bool empty(int trackIndex) const { return counts[trackIndex] == 0; } + int size(int trackIndex) const { return counts[trackIndex]; } + TrackExtensionCandidateN* begin(int trackIndex) { return candidates.data() + getFlatTrackExtensionCandidateIndex(trackIndex, 0); } + TrackExtensionCandidateN* end(int trackIndex) { return begin(trackIndex) + counts[trackIndex]; } + TrackExtensionCandidateN& get(int trackIndex, int candidateIndex) { return candidates[getFlatTrackExtensionCandidateIndex(trackIndex, candidateIndex)]; } + const TrackExtensionCandidateN& get(int trackIndex, int candidateIndex) const { return candidates[getFlatTrackExtensionCandidateIndex(trackIndex, candidateIndex)]; } + TrackExtensionCandidateN& getFlat(int flatIndex) { return candidates[flatIndex]; } + + std::vector candidates; + std::vector counts; + }; + + struct TrackFollowerScratch { + std::vector> activeHypotheses; + std::vector> nextHypotheses; + }; + + bool trackFollowing(TrackITSExt* track, bool outward, const int iteration, TrackFollowerScratch& scratch); + bool refitExtendedTrack(TrackITSExt& track, const int iteration); + void updateExtendedTrackTimeStamp(TrackITSExt& track, const int iteration); + virtual bool materializeTrackExtensionCandidate(TrackITSExt& track, const TrackExtensionCandidateN& candidate, const int iteration); + virtual void buildTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack); + void applyTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack); + o2::gpu::GPUChainITS* mChain = nullptr; TimeFrame* mTimeFrame; std::vector mTrkParams; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h index 69aa3c5fdaf06..054972faa8ed8 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h @@ -96,6 +96,10 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper 1) { + str += std::format(" Beam:{}", TrackFollowerBeamWidth); + } + } if (std::numeric_limits::max() != MaxMemory) { str += std::format(" MemLimit {:.2f} GB", double(MaxMemory) / constants::GB); } @@ -207,6 +217,15 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode p.RepeatRefitOut = tc.repeatRefitOut; p.ShiftRefToCluster = tc.shiftRefToCluster; p.CreateArtefactLabels = tc.createArtefactLabels; + p.TrackFollowerNSigmaCutZ = tc.trackFollowerNSigmaCutZ; + p.TrackFollowerNSigmaCutPhi = tc.trackFollowerNSigmaCutPhi; + p.TrackFollowerBeamWidth = std::max(1, tc.trackFollowerBeamWidth); + if (tc.trackFollower & 0x1) { + p.PassFlags.set(IterationStep::TrackFollowerTop); + } + if (tc.trackFollower & 0x2) { + p.PassFlags.set(IterationStep::TrackFollowerBot); + } p.PrintMemory = tc.printMemory; p.MaxMemory = tc.maxMemory; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index 8375004cbfbad..6ea1d84468312 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -475,6 +475,7 @@ template void TimeFrame::wipe() { deepVectorClear(mTracks); + deepVectorClear(mFittedExtensionTracks); deepVectorClear(mTracklets); deepVectorClear(mCells); deepVectorClear(mCellsNeighbours); diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index f17d961fc7bb7..1a1b24cb8d78a 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -20,6 +20,7 @@ #include "ITStracking/TrackingConfigParam.h" #include +#include #include #include #include @@ -92,6 +93,18 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e logger(std::format(" - Neighbours finding: {} neighbours found in {:.2f} ms", nNeighbours, timeNeighbours)); logger(std::format(" - Track finding: {} tracks found in {:.2f} ms", nTracks + mTimeFrame->getNumberOfTracks(), timeRoads)); total += timeTracklets + timeCells + timeNeighbours + timeRoads; + if (mTraits->supportsExtendTracks() && (mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot])) { + const int nClustersBefore = mTimeFrame->getNumberOfUsedClusters(); + const int nTracksBefore = std::count_if(mTimeFrame->getTracks().begin(), mTimeFrame->getTracks().end(), [](const auto& track) { + return track.getPattern() & 0xff000000; + }); + const auto timeExtending = evaluateTask(&Tracker::extendTracks, StateNames[mCurStep = Extending], iteration, evalLog, iteration); + const int nTracksAfter = std::count_if(mTimeFrame->getTracks().begin(), mTimeFrame->getTracks().end(), [](const auto& track) { + return track.getPattern() & 0xff000000; + }); + total += timeExtending; + logger(std::format(" - Extending tracks: {} tracks using {} clusters in {:.2f} ms", nTracksAfter - nTracksBefore, mTimeFrame->getNumberOfUsedClusters() - nClustersBefore, timeExtending)); + } } } catch (const BoundedMemoryResource::MemoryLimitExceeded& err) { handleException(err); diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index c4439dc74d29e..39c834ca3ec55 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -14,9 +14,14 @@ /// #include +#include #include +#include +#include +#include #include #include +#include #include #include @@ -30,6 +35,7 @@ #include "ITStracking/LayerMask.h" #include "ITStracking/ROFLookupTables.h" #include "ITStracking/TrackerTraits.h" +#include "ITStracking/TrackFollower.h" #include "ITStracking/TrackHelpers.h" #include "ITStracking/Tracklet.h" @@ -907,6 +913,378 @@ void TrackerTraits::markTracks(int iteration) } } +template +void TrackerTraits::extendTracks(const int iteration) +{ + const auto nTracks = mTimeFrame->getTracks().size(); + TrackExtensionCandidates candidatesPerTrack(nTracks); + mTimeFrame->mFittedExtensionTracks.clear(); + buildTrackExtensionCandidates(iteration, candidatesPerTrack); + applyTrackExtensionCandidates(iteration, candidatesPerTrack); + mTimeFrame->mFittedExtensionTracks.clear(); +} + +template +bool TrackerTraits::refitExtendedTrack(TrackITSExt& track, const int iteration) +{ + const auto propagator = o2::base::Propagator::Instance(); + const TrackingFrameInfo* tfInfos[NLayers]{}; + for (int iLayer = 0; iLayer < NLayers; ++iLayer) { + tfInfos[iLayer] = mTimeFrame->getTrackingFrameInfoOnLayer(iLayer).data(); + } + + o2::track::TrackPar linRef{track}; + track::resetTrackCovariance(track); + track.setChi2(0); + bool fitSuccess = track::fitTrack(track, + 0, + mTrkParams[iteration].NLayers, + 1, + mTrkParams[iteration].MaxChi2ClusterAttachment, + mTrkParams[iteration].MaxChi2NDF, + o2::constants::math::VeryBig, + 0, + mBz, + tfInfos, + mTrkParams[iteration].LayerxX0.data(), + propagator, + mTrkParams[iteration].CorrType, + &linRef, + mTrkParams[iteration].ShiftRefToCluster); + if (!fitSuccess) { + return false; + } + + track.getParamOut() = track.getParamIn(); + linRef = track.getParamOut(); + track::resetTrackCovariance(track); + track.setChi2(0); + fitSuccess = track::fitTrack(track, + mTrkParams[iteration].NLayers - 1, + -1, + -1, + mTrkParams[iteration].MaxChi2ClusterAttachment, + mTrkParams[iteration].MaxChi2NDF, + 50.f, + 0, + mBz, + tfInfos, + mTrkParams[iteration].LayerxX0.data(), + propagator, + mTrkParams[iteration].CorrType, + &linRef, + mTrkParams[iteration].ShiftRefToCluster); + return fitSuccess; +} + +template +void TrackerTraits::updateExtendedTrackTimeStamp(TrackITSExt& track, const int iteration) +{ + bool firstCluster{true}, nominalCompatible{true}; + TimeEstBC nominalTS, expandedTS; + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + const int cluster = track.getClusterIndex(iLayer); + if (cluster == constants::UnusedIndex) { + continue; + } + const int rof = mTimeFrame->getClusterROF(iLayer, cluster); + const auto nominalROFTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(rof); + const auto expandedROFTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(rof, true); + if (firstCluster) { + firstCluster = false; + nominalTS = nominalROFTS; + expandedTS = expandedROFTS; + continue; + } + if (nominalCompatible) { + if (nominalTS.isCompatible(nominalROFTS)) { + nominalTS += nominalROFTS; + } else { + nominalCompatible = false; + } + } + if (!expandedTS.isCompatible(expandedROFTS)) { + LOGP(fatal, "Clusters of an accepted track have non-overlapping expanded ROF time windows: {}+/-{} vs {}+/-{}", expandedROFTS.getTimeStamp(), expandedROFTS.getTimeStampError(), expandedTS.getTimeStamp(), expandedTS.getTimeStampError()); + } + expandedTS += expandedROFTS; + } + track.getTimeStamp() = (nominalCompatible ? nominalTS : expandedTS).makeSymmetrical(); +} + +template +bool TrackerTraits::materializeTrackExtensionCandidate(TrackITSExt& track, const TrackExtensionCandidateN& candidate, const int /*iteration*/) +{ + if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast(mTimeFrame->mFittedExtensionTracks.size())) { + return false; + } + track = mTimeFrame->mFittedExtensionTracks[candidate.resultIndex]; + return true; +} + +template +void TrackerTraits::buildTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack) +{ + struct ThreadExtensionResults { + std::vector tracks; + std::vector candidateIndicesToPatch; + }; + tbb::enumerable_thread_specific fittedTracks; + + auto prepareCandidate = [&](int trackIndex, const TrackITSExt& backup, TrackITSExt& candidate) { + if (!refitExtendedTrack(candidate, iteration)) { + return; + } + updateExtendedTrackTimeStamp(candidate, iteration); + const auto diff = (candidate.getPattern() & ~backup.getPattern()) & makeAddedClustersPatternMask(); + if (!diff) { + return; + } + applyExtendedClustersPattern(candidate, diff); + + TrackExtensionCandidateN extension; + extension.trackIndex = trackIndex; + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + if (backup.getClusterIndex(iLayer) == constants::UnusedIndex && candidate.getClusterIndex(iLayer) != constants::UnusedIndex) { + extension.addedClusters[iLayer] = candidate.getClusterIndex(iLayer); + ++extension.nAddedClusters; + } + } + if (!extension.nAddedClusters) { + return; + } + extension.chi2 = candidate.getChi2(); + const int candidateIndex = candidatesPerTrack.add(trackIndex, extension); + if (candidateIndex < 0) { + return; + } + auto& storedExtension = candidatesPerTrack.getFlat(candidateIndex); + auto& localFittedTracks = fittedTracks.local(); + storedExtension.resultIndex = static_cast(localFittedTracks.tracks.size()); + localFittedTracks.tracks.push_back(candidate); + localFittedTracks.candidateIndicesToPatch.push_back(candidateIndex); + }; + + const bool extendTop = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; + const bool extendBot = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; + auto& tracks = mTimeFrame->getTracks(); + tbb::enumerable_thread_specific trackFollowerScratch; + + const uint32_t lastLayer = static_cast(mTrkParams[iteration].NLayers - 1); + + auto buildCandidates = [&](int iTrack) { + const auto& backup = tracks[iTrack]; + auto& scratch = trackFollowerScratch.local(); + + std::optional topResult, botResult; + + if (extendTop && backup.getLastClusterLayer() != lastLayer) { + auto candidate{backup}; + if (trackFollowing(&candidate, true, iteration, scratch)) { + topResult = candidate; + prepareCandidate(iTrack, backup, candidate); + } + } + if (extendBot && backup.getFirstClusterLayer() != 0) { + auto candidate{backup}; + if (trackFollowing(&candidate, false, iteration, scratch)) { + botResult = candidate; + prepareCandidate(iTrack, backup, candidate); + } + } + if (extendTop && extendBot) { + if (topResult && topResult->getFirstClusterLayer() != 0) { + auto candidate = *topResult; + if (trackFollowing(&candidate, false, iteration, scratch)) { + prepareCandidate(iTrack, backup, candidate); + } + } + if (botResult && botResult->getLastClusterLayer() != lastLayer) { + auto candidate = *botResult; + if (trackFollowing(&candidate, true, iteration, scratch)) { + prepareCandidate(iTrack, backup, candidate); + } + } + } + }; + + if (mTaskArena->max_concurrency() <= 1) { + for (int iTrack{0}; iTrack < static_cast(tracks.size()); ++iTrack) { + buildCandidates(iTrack); + } + } else { + mTaskArena->execute([&] { + tbb::parallel_for(0, static_cast(tracks.size()), buildCandidates); + }); + } + + size_t nFittedExtensionTracks{0}; + for (auto& localFittedTracks : fittedTracks) { + nFittedExtensionTracks += localFittedTracks.tracks.size(); + } + mTimeFrame->mFittedExtensionTracks.reserve(nFittedExtensionTracks); + + int resultOffset{0}; + for (auto& localFittedTracks : fittedTracks) { + for (auto candidateIndex : localFittedTracks.candidateIndicesToPatch) { + candidatesPerTrack.getFlat(candidateIndex).resultIndex += resultOffset; + } + mTimeFrame->mFittedExtensionTracks.insert(mTimeFrame->mFittedExtensionTracks.end(), localFittedTracks.tracks.begin(), localFittedTracks.tracks.end()); + resultOffset += static_cast(localFittedTracks.tracks.size()); + } +} + +template +void TrackerTraits::applyTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack) +{ + auto& tracks = mTimeFrame->getTracks(); + + for (int iTrack{0}; iTrack < static_cast(tracks.size()); ++iTrack) { + std::stable_sort(candidatesPerTrack.begin(iTrack), candidatesPerTrack.end(iTrack), isBetterTrackExtensionCandidate); + while (!candidatesPerTrack.empty(iTrack) && (candidatesPerTrack.get(iTrack, candidatesPerTrack.size(iTrack) - 1).nAddedClusters <= 0)) { + candidatesPerTrack.pop_back(iTrack); + } + } + + std::array, NLayers> claimedClusters; + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + claimedClusters[iLayer].resize(mTimeFrame->getClusters()[iLayer].size(), 0); + } + + struct Entry { + int track; + int idx; + }; + auto cmp = [&](const Entry& a, const Entry& b) { + const auto& ca = candidatesPerTrack.get(a.track, a.idx); + const auto& cb = candidatesPerTrack.get(b.track, b.idx); + if (isBetterTrackExtensionCandidate(cb, ca)) { + return true; + } + if (isBetterTrackExtensionCandidate(ca, cb)) { + return false; + } + if (a.track != b.track) { + return a.track > b.track; + } + return a.idx > b.idx; + }; + std::priority_queue, decltype(cmp)> pq(cmp); + for (int iTrack{0}; iTrack < static_cast(tracks.size()); ++iTrack) { + if (!candidatesPerTrack.empty(iTrack)) { + pq.push({iTrack, 0}); + } + } + + auto tryNext = [&](int trackIndex, int idx) { + if (idx + 1 < candidatesPerTrack.size(trackIndex)) { + pq.push({trackIndex, idx + 1}); + } + }; + + while (!pq.empty()) { + const Entry e = pq.top(); + pq.pop(); + const auto& candidate = candidatesPerTrack.get(e.track, e.idx); + + bool hasContention{false}; + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + const int cluster = candidate.addedClusters[iLayer]; + if (cluster == constants::UnusedIndex) { + continue; + } + if (cluster >= static_cast(claimedClusters[iLayer].size()) || claimedClusters[iLayer][cluster]) { + hasContention = true; + break; + } + } + if (hasContention) { + tryNext(e.track, e.idx); + continue; + } + auto extendedTrack = tracks[e.track]; + if (!materializeTrackExtensionCandidate(extendedTrack, candidate, iteration)) { + tryNext(e.track, e.idx); + continue; + } + tracks[e.track] = extendedTrack; + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + const int cluster = candidate.addedClusters[iLayer]; + if (cluster == constants::UnusedIndex) { + continue; + } + claimedClusters[iLayer][cluster] = 1; + mTimeFrame->markUsedCluster(iLayer, cluster); + } + } +} + +template +bool TrackerTraits::trackFollowing(TrackITSExt* track, bool outward, const int iteration, TrackFollowerScratch& scratch) +{ + const int beamWidth = std::max(1, mTrkParams[iteration].TrackFollowerBeamWidth); + if (static_cast(scratch.activeHypotheses.size()) < beamWidth) { + scratch.activeHypotheses.resize(beamWidth); + } + if (static_cast(scratch.nextHypotheses.size()) < beamWidth) { + scratch.nextHypotheses.resize(beamWidth); + } + + const Cluster* clustersPtrs[NLayers]{}; + const unsigned char* usedClustersPtrs[NLayers]{}; + const int* clustersIndexTablesPtrs[NLayers]{}; + const int* rofClustersPtrs[NLayers]{}; + const TrackingFrameInfo* tfInfoPtrs[NLayers]{}; + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + clustersPtrs[iLayer] = mTimeFrame->getClusters()[iLayer].data(); + usedClustersPtrs[iLayer] = mTimeFrame->getUsedClusters(iLayer).data(); + clustersIndexTablesPtrs[iLayer] = mTimeFrame->getIndexTable(0, iLayer).data(); + rofClustersPtrs[iLayer] = mTimeFrame->getROFrameClusters(iLayer).data(); + tfInfoPtrs[iLayer] = mTimeFrame->getTrackingFrameInfoOnLayer(iLayer).data(); + } + + TrackITSExt updated; + const bool ok = followTrackExtensionDirection( + *track, + mTimeFrame->getIndexTableUtils(), + mTimeFrame->getROFMaskView(), + mTimeFrame->getROFOverlapTableView(), + clustersPtrs, + usedClustersPtrs, + clustersIndexTablesPtrs, + rofClustersPtrs, + tfInfoPtrs, + mTrkParams[iteration].LayerRadii.data(), + mTrkParams[iteration].LayerxX0.data(), + mTrkParams[iteration].NLayers, + mTrkParams[iteration].PhiBins, + beamWidth, + mBz, + mTrkParams[iteration].MaxChi2ClusterAttachment, + mTrkParams[iteration].MaxChi2NDF, + mTrkParams[iteration].TrackFollowerNSigmaCutPhi, + mTrkParams[iteration].TrackFollowerNSigmaCutZ, + outward, + o2::base::Propagator::Instance(), + mTrkParams[iteration].CorrType, + scratch.activeHypotheses.data(), + scratch.nextHypotheses.data(), + updated); + if (!ok) { + return false; + } + + auto& trackParam = outward ? track->getParamOut() : track->getParamIn(); + trackParam = outward ? updated.getParamOut() : updated.getParamIn(); + track->setChi2(updated.getChi2()); + track->getTimeStamp() = updated.getTimeStamp(); + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + if (track->getClusterIndex(iLayer) == constants::UnusedIndex && updated.getClusterIndex(iLayer) != constants::UnusedIndex) { + track->setExternalClusterIndex(iLayer, updated.getClusterIndex(iLayer), true); + } + } + return true; +} + template void TrackerTraits::setBz(float bz) { From d4e7f5a020e97103f9db4a15516d9ce79af9bd1c Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Wed, 20 May 2026 18:54:00 +0200 Subject: [PATCH 2/4] ITS: integrate track extension in road finding --- .../GPU/ITStrackingGPU/TrackerTraitsGPU.h | 6 - .../GPU/ITStrackingGPU/TrackingKernels.h | 18 +- .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 154 ++---- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 370 ++++++++++++++- .../tracking/include/ITStracking/TimeFrame.h | 1 - .../ITStracking/TrackExtensionCandidate.h | 22 +- .../tracking/include/ITStracking/Tracker.h | 1 - .../include/ITStracking/TrackerTraits.h | 52 +-- .../include/ITStracking/TrackingConfigParam.h | 2 +- .../ITSMFT/ITS/tracking/src/TimeFrame.cxx | 1 - Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx | 12 - .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 442 +++++------------- 12 files changed, 551 insertions(+), 530 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h index 13773ac234027..81d870c5b46c2 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h @@ -35,7 +35,6 @@ class TrackerTraitsGPU final : public TrackerTraits void computeLayerCells(const int iteration) final; void findCellsNeighbours(const int iteration) final; void findRoads(const int iteration) final; - void extendTracks(const int iteration) final; void setBz(float) final; @@ -48,11 +47,6 @@ class TrackerTraitsGPU final : public TrackerTraits int getTFNumberOfCells() const override; private: - bool hasTrackFollower(const int iteration) const; - - void buildTrackExtensionCandidates(const int iteration, typename TrackerTraits::TrackExtensionCandidates& candidatesPerTrack) final; - bool materializeTrackExtensionCandidate(TrackITSExt& track, const typename TrackerTraits::TrackExtensionCandidateN& candidate, const int iteration) final; - IndexTableUtilsN* mDeviceIndexTableUtils; gpu::TimeFrameGPU* mTimeFrameGPU; }; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 3e50aedab5323..ff541e0e5a839 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -262,7 +262,6 @@ void countTrackSeedHandler(TrackSeed* trackSeeds, const std::vector& layerxX0Host, const unsigned int nSeeds, const float Bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, @@ -276,20 +275,35 @@ template void computeTrackSeedHandler(TrackSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, const Cluster** unsortedClusters, + const IndexTableUtils* utils, + const typename ROFMaskTable::View& rofMask, + const typename ROFOverlapTable::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, o2::its::TrackITSExt* tracks, const int* seedLUT, + TrackExtensionHypothesis* activeHypotheses, + TrackExtensionHypothesis* nextHypotheses, const std::vector& layerRadiiHost, const std::vector& minPtsHost, const std::vector& layerxX0Host, const unsigned int nSeeds, const unsigned int nTracks, const float Bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float nSigmaCutPhi, + const float nSigmaCutZ, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, o2::its::ExternalAllocator* alloc); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index b88b63d04e053..43c45649b656a 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -304,10 +304,13 @@ template void TrackerTraitsGPU::findRoads(const int iteration) { bounded_vector> firstClusters(this->mTrkParams[iteration].NLayers, bounded_vector(this->getMemoryPool().get()), this->getMemoryPool().get()); - bounded_vector> sharedFirstClusters(this->mTrkParams[iteration].NLayers, bounded_vector(this->getMemoryPool().get()), this->getMemoryPool().get()); firstClusters.resize(this->mTrkParams[iteration].NLayers); - sharedFirstClusters.resize(this->mTrkParams[iteration].NLayers); const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView(); + const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; + const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; + const bool extendTracks = extendTop || extendBot; + size_t nExtendedTracks{0}; + size_t nExtendedClusters{0}; 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) { @@ -356,7 +359,6 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].LayerxX0, trackSeeds.size(), this->mBz, - startLevel, this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].MaxChi2NDF, this->mTrkParams[iteration].ReseedIfShorter, @@ -366,153 +368,57 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].CorrType, mTimeFrameGPU->getFrameworkAllocator()); mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size()); + if (extendTracks) { + mTimeFrameGPU->createTrackExtensionScratchDevice(kTrackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerBeamWidth); + } computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), mTimeFrameGPU->getDeviceArrayUnsortedClusters(), + mTimeFrameGPU->getDeviceIndexTableUtils(), + mTimeFrameGPU->getDeviceROFMaskTableView(), + mTimeFrameGPU->getDeviceROFOverlapTableView(), + mTimeFrameGPU->getDeviceArrayClusters(), + (const unsigned char**)mTimeFrameGPU->getDeviceArrayUsedClusters(), + mTimeFrameGPU->getDeviceArrayClustersIndexTables(), + mTimeFrameGPU->getDeviceROFrameClusters(), mTimeFrameGPU->getDeviceTrackITSExt(), mTimeFrameGPU->getDeviceTrackSeedsLUT(), + extendTracks ? mTimeFrameGPU->getDeviceActiveTrackExtensionHypotheses() : nullptr, + extendTracks ? mTimeFrameGPU->getDeviceNextTrackExtensionHypotheses() : nullptr, this->mTrkParams[iteration].LayerRadii, this->mTrkParams[iteration].MinPt, this->mTrkParams[iteration].LayerxX0, trackSeeds.size(), mTimeFrameGPU->getNTrackSeeds(), this->mBz, - startLevel, this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].MaxChi2NDF, this->mTrkParams[iteration].ReseedIfShorter, this->mTrkParams[iteration].RepeatRefitOut, this->mTrkParams[iteration].ShiftRefToCluster, + this->mTrkParams[iteration].NLayers, + this->mTrkParams[iteration].PhiBins, + this->mTrkParams[iteration].TrackFollowerBeamWidth, + extendTop, + extendBot, + this->mTrkParams[iteration].TrackFollowerNSigmaCutPhi, + this->mTrkParams[iteration].TrackFollowerNSigmaCutZ, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[iteration].CorrType, mTimeFrameGPU->getFrameworkAllocator()); mTimeFrameGPU->downloadTrackITSExtDevice(); auto& tracks = mTimeFrameGPU->getTrackITSExt(); - this->acceptTracks(iteration, tracks, firstClusters); + this->acceptTracks(iteration, tracks, firstClusters, nExtendedTracks, nExtendedClusters); mTimeFrameGPU->loadUsedClustersDevice(); } - this->markTracks(iteration); - if (!hasTrackFollower(iteration)) { - // wipe the artefact memory - mTimeFrameGPU->popMemoryStack(iteration); + if (extendTracks) { + LOGP(info, "Integrated track extension accepted {} tracks using {} clusters in iteration {}", nExtendedTracks, nExtendedClusters, iteration); } -}; - -template -void TrackerTraitsGPU::extendTracks(const int iteration) -{ - TrackerTraits::extendTracks(iteration); - mTimeFrameGPU->loadUsedClustersDevice(); - // wipe the artefact memory kept alive for the track follower + this->markTracks(iteration); + // wipe the artefact memory mTimeFrameGPU->popMemoryStack(iteration); -} - -template -bool TrackerTraitsGPU::hasTrackFollower(const int iteration) const -{ - return this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || - this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; -} - -template -void TrackerTraitsGPU::buildTrackExtensionCandidates(const int iteration, typename TrackerTraits::TrackExtensionCandidates& candidatesPerTrack) -{ - const auto nTracks = this->mTimeFrame->getTracks().size(); - const int beamWidth = std::max(1, this->mTrkParams[iteration].TrackFollowerBeamWidth); - mTimeFrameGPU->syncStreams(); - mTimeFrameGPU->loadTrackExtensionStartTracksDevice(); - mTimeFrameGPU->createTrackExtensionCandidatesDevice(nTracks); - mTimeFrameGPU->createTrackExtensionScratchDevice(kTrackExtensionLaunchThreads, beamWidth); - std::array layerRadii{}; - std::array layerxX0{}; - for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) { - layerRadii[iLayer] = this->mTrkParams[iteration].LayerRadii[iLayer]; - layerxX0[iLayer] = this->mTrkParams[iteration].LayerxX0[iLayer]; - } - computeTrackExtensionCandidatesHandler(mTimeFrameGPU->getDeviceTrackExtensionStartTracks(), - mTimeFrameGPU->getDeviceIndexTableUtils(), - mTimeFrameGPU->getDeviceROFMaskTableView(), - mTimeFrameGPU->getDeviceROFOverlapTableView(), - mTimeFrameGPU->getDeviceArrayClusters(), - (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(), - mTimeFrameGPU->getDeviceArrayClustersIndexTables(), - mTimeFrameGPU->getDeviceROFrameClusters(), - mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), - mTimeFrameGPU->getDeviceTrackExtensionCandidates(), - mTimeFrameGPU->getDeviceTrackExtensionCandidateOffsets(), - mTimeFrameGPU->getDeviceActiveTrackExtensionHypotheses(), - mTimeFrameGPU->getDeviceNextTrackExtensionHypotheses(), - layerRadii, - layerxX0, - static_cast(nTracks), - this->mTrkParams[iteration].NLayers, - this->mTrkParams[iteration].PhiBins, - beamWidth, - this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop], - this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot], - this->mBz, - this->mTrkParams[iteration].MaxChi2ClusterAttachment, - this->mTrkParams[iteration].MaxChi2NDF, - this->mTrkParams[iteration].TrackFollowerNSigmaCutPhi, - this->mTrkParams[iteration].TrackFollowerNSigmaCutZ, - mTimeFrameGPU->getDevicePropagator(), - this->mTrkParams[iteration].CorrType, - mTimeFrameGPU->getStream(0)); - mTimeFrameGPU->createTrackExtensionResultsDevice(nTracks); - computeTrackExtensionResultsHandler(mTimeFrameGPU->getDeviceTrackExtensionStartTracks(), - mTimeFrameGPU->getDeviceTrackExtensionCandidates(), - mTimeFrameGPU->getDeviceTrackExtensionCandidateOffsets(), - mTimeFrameGPU->getDeviceTrackExtensionResults(), - mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), - layerxX0, - static_cast(nTracks), - this->mTrkParams[iteration].NLayers, - this->mBz, - this->mTrkParams[iteration].MaxChi2ClusterAttachment, - this->mTrkParams[iteration].MaxChi2NDF, - mTimeFrameGPU->getDevicePropagator(), - this->mTrkParams[iteration].CorrType, - this->mTrkParams[iteration].ShiftRefToCluster, - mTimeFrameGPU->getStream(0)); - mTimeFrameGPU->downloadTrackExtensionResultsDevice(); - - const auto& results = mTimeFrameGPU->getTrackExtensionResults(); - for (int iResult{0}; iResult < static_cast(results.size()); ++iResult) { - const auto& result = results[iResult]; - if (!result.isValid()) { - continue; - } - auto candidate = result.candidate; - candidate.resultIndex = iResult; - if (candidatesPerTrack.add(candidate.trackIndex, candidate) < 0) { - continue; - } - } -} - -template -bool TrackerTraitsGPU::materializeTrackExtensionCandidate(TrackITSExt& track, const typename TrackerTraits::TrackExtensionCandidateN& candidate, const int iteration) -{ - const auto& results = mTimeFrameGPU->getTrackExtensionResults(); - if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast(results.size())) { - return TrackerTraits::materializeTrackExtensionCandidate(track, candidate, iteration); - } - const auto& result = results[candidate.resultIndex]; - if (!result.isValid() || result.candidate.trackIndex != candidate.trackIndex) { - return false; - } - track = result.track; - this->updateExtendedTrackTimeStamp(track, iteration); - uint32_t diff{0}; - for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) { - if (candidate.addedClusters[iLayer] != constants::UnusedIndex) { - diff |= (0x1u << iLayer); - } - } - applyExtendedClustersPattern(track, diff); - return true; -} +}; template int TrackerTraitsGPU::getTFNumberOfClusters() const diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index d42b373ca3e64..55a0bc4d069e0 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -285,6 +285,85 @@ GPUdi() bool fitTrackExtensionResult(const TrackITSExt& startTrack, return true; } +template +GPUdi() bool refitTrackExtensionResult(TrackITSExt& track, + const TrackingFrameInfo* const* trackingFrameInfo, + const float* layerxX0, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster) +{ + o2::track::TrackPar linRef{track}; + o2::its::track::resetTrackCovariance(track); + track.setChi2(0); + bool fitSuccess = o2::its::track::fitTrack(track, + 0, + nLayers, + 1, + maxChi2ClusterAttachment, + maxChi2NDF, + o2::constants::math::VeryBig, + 0, + bz, + trackingFrameInfo, + layerxX0, + propagator, + matCorrType, + &linRef, + shiftRefToCluster); + if (!fitSuccess) { + return false; + } + + track.getParamOut() = track.getParamIn(); + linRef = track.getParamOut(); + o2::its::track::resetTrackCovariance(track); + track.setChi2(0); + return o2::its::track::fitTrack(track, + nLayers - 1, + -1, + -1, + maxChi2ClusterAttachment, + maxChi2NDF, + 50.f, + 0, + bz, + trackingFrameInfo, + layerxX0, + propagator, + matCorrType, + &linRef, + shiftRefToCluster); +} + +template +GPUdi() void finaliseTrackExtensionCandidate(const uint32_t backupPattern, + TrackITSExt& candidate, + const TrackingFrameInfo* const* trackingFrameInfo, + const float* layerxX0, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + TrackITSExt& best) +{ + const auto diff = (candidate.getPattern() & ~backupPattern) & makeAddedClustersPatternMask(); + if (!diff || !refitTrackExtensionResult(candidate, trackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster)) { + return; + } + applyExtendedClustersPattern(candidate, diff); + if (o2::its::track::isBetter(candidate, best)) { + best = candidate; + } +} + template GPUg() void __launch_bounds__(256, 1) computeTrackExtensionResultsKernel(const TrackITSExt* tracks, const TrackExtensionCandidate* candidates, @@ -333,33 +412,86 @@ GPUg() void __launch_bounds__(256, 1) computeTrackExtensionResultsKernel(const T } } -template +template +GPUg() void __launch_bounds__(256, 1) countTrackSeedsKernel( + TrackSeed* trackSeeds, + const TrackingFrameInfo** foundTrackingFrameInfo, + const Cluster** unsortedClusters, + int* seedLUT, + const float* layerRadii, + const float* minPts, + const float* layerxX0, + const unsigned int nSeeds, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const int reseedIfShorter, + const bool repeatRefitOut, + const bool shiftRefToCluster, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType) +{ + for (int iCurrentTrackSeedIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackSeedIndex < nSeeds; iCurrentTrackSeedIndex += blockDim.x * gridDim.x) { + TrackITSExt temporaryTrack; + if (o2::its::track::refitTrack(trackSeeds[iCurrentTrackSeedIndex], + temporaryTrack, + maxChi2ClusterAttachment, + maxChi2NDF, + bz, + foundTrackingFrameInfo, + unsortedClusters, + layerxX0, + layerRadii, + minPts, + propagator, + matCorrType, + reseedIfShorter, + shiftRefToCluster, + repeatRefitOut)) { + seedLUT[iCurrentTrackSeedIndex] = 1; + } + } +} + +template GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( TrackSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, const Cluster** unsortedClusters, + const IndexTableUtils* utils, + const typename ROFMaskTable::View rofMask, + const typename ROFOverlapTable::View rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, o2::its::TrackITSExt* tracks, - maybe_const* seedLUT, + const int* seedLUT, + TrackExtensionHypothesis* activeHypothesesScratch, + TrackExtensionHypothesis* nextHypothesesScratch, const float* layerRadii, const float* minPts, const float* layerxX0, const unsigned int nSeeds, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, + const int nLayers, + const int phiBins, + const int beamWidthConfig, + const bool extendTop, + const bool extendBot, + const float nSigmaCutPhi, + const float nSigmaCutZ, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType) { for (int iCurrentTrackSeedIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackSeedIndex < nSeeds; iCurrentTrackSeedIndex += blockDim.x * gridDim.x) { - - if constexpr (!initRun) { - if (seedLUT[iCurrentTrackSeedIndex] == seedLUT[iCurrentTrackSeedIndex + 1]) { - continue; - } + if (seedLUT[iCurrentTrackSeedIndex] == seedLUT[iCurrentTrackSeedIndex + 1]) { + continue; } TrackITSExt temporaryTrack; bool refitSuccess = o2::its::track::refitTrack(trackSeeds[iCurrentTrackSeedIndex], @@ -378,11 +510,148 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( shiftRefToCluster, repeatRefitOut); if (refitSuccess) { - if constexpr (initRun) { - seedLUT[iCurrentTrackSeedIndex] = 1; - } else { - tracks[seedLUT[iCurrentTrackSeedIndex]] = temporaryTrack; + if ((extendTop || extendBot) && activeHypothesesScratch && nextHypothesesScratch) { + const int beamWidth = o2::gpu::CAMath::Max(beamWidthConfig, 1); + const int threadIndex = blockIdx.x * blockDim.x + threadIdx.x; + auto* activeHypotheses = activeHypothesesScratch + threadIndex * beamWidth; + auto* nextHypotheses = nextHypothesesScratch + threadIndex * beamWidth; + const auto backupPattern = temporaryTrack.getPattern(); + auto best = temporaryTrack; + TrackITSExt topResult; + TrackITSExt botResult; + bool hasTopResult{false}; + bool hasBotResult{false}; + const uint32_t lastLayer = static_cast(nLayers - 1); + + if (extendTop && getTrackExtensionLastClusterLayer(temporaryTrack) != lastLayer) { + auto candidate = temporaryTrack; + if (followTrackExtensionDirection(temporaryTrack, + *utils, + rofMask, + rofOverlaps, + clusters, + usedClusters, + clustersIndexTables, + ROFClusters, + foundTrackingFrameInfo, + layerRadii, + layerxX0, + nLayers, + phiBins, + beamWidth, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + nSigmaCutPhi, + nSigmaCutZ, + true, + propagator, + matCorrType, + activeHypotheses, + nextHypotheses, + candidate)) { + topResult = candidate; + hasTopResult = true; + finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + } + } + if (extendBot && getTrackExtensionFirstClusterLayer(temporaryTrack) != 0) { + auto candidate = temporaryTrack; + if (followTrackExtensionDirection(temporaryTrack, + *utils, + rofMask, + rofOverlaps, + clusters, + usedClusters, + clustersIndexTables, + ROFClusters, + foundTrackingFrameInfo, + layerRadii, + layerxX0, + nLayers, + phiBins, + beamWidth, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + nSigmaCutPhi, + nSigmaCutZ, + false, + propagator, + matCorrType, + activeHypotheses, + nextHypotheses, + candidate)) { + botResult = candidate; + hasBotResult = true; + finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + } + } + if (extendTop && extendBot) { + if (hasTopResult && getTrackExtensionFirstClusterLayer(topResult) != 0) { + auto candidate = topResult; + if (followTrackExtensionDirection(topResult, + *utils, + rofMask, + rofOverlaps, + clusters, + usedClusters, + clustersIndexTables, + ROFClusters, + foundTrackingFrameInfo, + layerRadii, + layerxX0, + nLayers, + phiBins, + beamWidth, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + nSigmaCutPhi, + nSigmaCutZ, + false, + propagator, + matCorrType, + activeHypotheses, + nextHypotheses, + candidate)) { + finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + } + } + if (hasBotResult && getTrackExtensionLastClusterLayer(botResult) != lastLayer) { + auto candidate = botResult; + if (followTrackExtensionDirection(botResult, + *utils, + rofMask, + rofOverlaps, + clusters, + usedClusters, + clustersIndexTables, + ROFClusters, + foundTrackingFrameInfo, + layerRadii, + layerxX0, + nLayers, + phiBins, + beamWidth, + bz, + maxChi2ClusterAttachment, + maxChi2NDF, + nSigmaCutPhi, + nSigmaCutZ, + true, + propagator, + matCorrType, + activeHypotheses, + nextHypotheses, + candidate)) { + finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + } + } + } + temporaryTrack = best; } + tracks[seedLUT[iCurrentTrackSeedIndex]] = temporaryTrack; } } } @@ -1375,7 +1644,6 @@ void countTrackSeedHandler(TrackSeed* trackSeeds, const std::vector& layerxX0Host, const unsigned int nSeeds, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, @@ -1391,18 +1659,16 @@ void countTrackSeedHandler(TrackSeed* trackSeeds, thrust::device_vector minPts(minPtsHost); thrust::device_vector layerRadii(layerRadiiHost); thrust::device_vector layerxX0(layerxX0Host); - gpu::fitTrackSeedsKernel<<<60, 256>>>( + gpu::countTrackSeedsKernel<<<60, 256>>>( trackSeeds, // CellSeed* foundTrackingFrameInfo, // TrackingFrameInfo** unsortedClusters, // Cluster** - nullptr, // TrackITSExt* seedLUT, // int* thrust::raw_pointer_cast(&layerRadii[0]), // const float* thrust::raw_pointer_cast(&minPts[0]), // const float* thrust::raw_pointer_cast(&layerxX0[0]), // const float* nSeeds, // const unsigned int bz, // const float - startLevel, // const int maxChi2ClusterAttachment, // float maxChi2NDF, // float reseedIfShorter, // int @@ -1418,20 +1684,35 @@ template void computeTrackSeedHandler(TrackSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, const Cluster** unsortedClusters, + const IndexTableUtils* utils, + const typename ROFMaskTable::View& rofMask, + const typename ROFOverlapTable::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, o2::its::TrackITSExt* tracks, const int* seedLUT, + TrackExtensionHypothesis* activeHypotheses, + TrackExtensionHypothesis* nextHypotheses, const std::vector& layerRadiiHost, const std::vector& minPtsHost, const std::vector& layerxX0Host, const unsigned int nSeeds, const unsigned int nTracks, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float nSigmaCutPhi, + const float nSigmaCutZ, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, o2::its::ExternalAllocator* alloc) @@ -1439,23 +1720,38 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, thrust::device_vector minPts(minPtsHost); thrust::device_vector layerRadii(layerRadiiHost); thrust::device_vector layerxX0(layerxX0Host); - gpu::fitTrackSeedsKernel<<<60, 256>>>( + gpu::fitTrackSeedsKernel<<<60, 256>>>( trackSeeds, // CellSeed* foundTrackingFrameInfo, // TrackingFrameInfo** unsortedClusters, // Cluster** + utils, // IndexTableUtils* + rofMask, // ROFMaskTable::View + rofOverlaps, // ROFOverlapTable::View + clusters, // Cluster** + usedClusters, // unsigned char** + clustersIndexTables, // int** + ROFClusters, // int** tracks, // TrackITSExt* seedLUT, // const int* + activeHypotheses, // TrackExtensionHypothesis* + nextHypotheses, // TrackExtensionHypothesis* thrust::raw_pointer_cast(&layerRadii[0]), // const float* thrust::raw_pointer_cast(&minPts[0]), // const float* thrust::raw_pointer_cast(&layerxX0[0]), // const float* nSeeds, // const unsigned int bz, // const float - startLevel, // const int maxChi2ClusterAttachment, // float maxChi2NDF, // float reseedIfShorter, // int repeatRefitOut, // bool shiftRefToCluster, // bool + nLayers, // int + phiBins, // int + beamWidth, // int + extendTop, // bool + extendBot, // bool + nSigmaCutPhi, // float + nSigmaCutZ, // float propagator, // const o2::base::Propagator* matCorrType); // o2::base::PropagatorF::MatCorrType auto sync_policy = THRUST_NAMESPACE::par(gpu::TypedAllocator(alloc)); @@ -1663,7 +1959,6 @@ template void countTrackSeedHandler(TrackSeed<7>* trackSeeds, const std::vector& layerxX0Host, const unsigned int nSeeds, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, @@ -1676,20 +1971,35 @@ template void countTrackSeedHandler(TrackSeed<7>* trackSeeds, template void computeTrackSeedHandler(TrackSeed<7>* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, const Cluster** unsortedClusters, + const IndexTableUtils<7>* utils, + const ROFMaskTable<7>::View& rofMask, + const ROFOverlapTable<7>::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, o2::its::TrackITSExt* tracks, const int* seedLUT, + TrackExtensionHypothesis<7>* activeHypotheses, + TrackExtensionHypothesis<7>* nextHypotheses, const std::vector& layerRadiiHost, const std::vector& minPtsHost, const std::vector& layerxX0Host, const unsigned int nSeeds, const unsigned int nTracks, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float nSigmaCutPhi, + const float nSigmaCutZ, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, o2::its::ExternalAllocator* alloc); @@ -1895,7 +2205,6 @@ template void countTrackSeedHandler(TrackSeed<11>* trackSeeds, const std::vector& layerxX0Host, const unsigned int nSeeds, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, @@ -1908,20 +2217,35 @@ template void countTrackSeedHandler(TrackSeed<11>* trackSeeds, template void computeTrackSeedHandler(TrackSeed<11>* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, const Cluster** unsortedClusters, + const IndexTableUtils<11>* utils, + const ROFMaskTable<11>::View& rofMask, + const ROFOverlapTable<11>::View& rofOverlaps, + const Cluster** clusters, + const unsigned char** usedClusters, + const int** clustersIndexTables, + const int** ROFClusters, o2::its::TrackITSExt* tracks, const int* seedLUT, + TrackExtensionHypothesis<11>* activeHypotheses, + TrackExtensionHypothesis<11>* nextHypotheses, const std::vector& layerRadiiHost, const std::vector& minPtsHost, const std::vector& layerxX0Host, const unsigned int nSeeds, const unsigned int nTracks, const float bz, - const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, + const int nLayers, + const int phiBins, + const int beamWidth, + const bool extendTop, + const bool extendBot, + const float nSigmaCutPhi, + const float nSigmaCutZ, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, o2::its::ExternalAllocator* alloc); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index 7d908d6265660..3fef2dc640cbc 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -279,7 +279,6 @@ struct TimeFrame { std::vector> mTracklets; std::vector> mCells; bounded_vector mTracks; - std::vector mFittedExtensionTracks; bounded_vector mTracksLabel; std::vector> mCellsNeighbours; std::vector> mCellsNeighboursTopology; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h index be165b54df8c6..5ff5bc4c0828b 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h @@ -36,14 +36,34 @@ GPUhdi() constexpr uint32_t makeAddedClustersPatternMask() template GPUhdi() void applyExtendedClustersPattern(TrackITSExt& track, uint32_t diff) { + diff &= makeAddedClustersPatternMask(); + track.setUserField(static_cast(diff)); if constexpr (NLayers <= kMaxLayersInTrackPattern) { track.setPattern(track.getPattern() | (diff << kExtendedPatternShift)); } else { (void)track; - (void)diff; } } +template +GPUhdi() uint32_t getAddedClustersPattern(const TrackITSExt& track) +{ + const auto mask = makeAddedClustersPatternMask(); + if constexpr (NLayers <= kMaxLayersInTrackPattern) { + const auto diff = (track.getPattern() >> kExtendedPatternShift) & mask; + if (diff) { + return diff; + } + } + return track.getUserField() & mask; +} + +GPUhdi() void clearAddedClustersPattern(TrackITSExt& track) +{ + track.setUserField(0); + track.getParamOut().setUserField(0); +} + template struct TrackExtensionHypothesis { o2::track::TrackParCov param; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h index daa185c945560..2362b6b2d9816 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h @@ -74,7 +74,6 @@ class Tracker void computeCells(int iteration) { mTraits->computeLayerCells(iteration); } void findCellsNeighbours(int iteration) { mTraits->findCellsNeighbours(iteration); } void findRoads(int iteration) { mTraits->findRoads(iteration); } - void extendTracks(int iteration) { mTraits->extendTracks(iteration); } void rectifyClusterIndices(); void sortTracks(); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index e870af0ad7baa..201ee0470d20b 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -20,6 +20,7 @@ #include #include +#include "DetectorsBase/Propagator.h" #include "ITStracking/Configuration.h" #include "ITStracking/Constants.h" #include "ITStracking/IndexTableUtils.h" @@ -57,13 +58,11 @@ class TrackerTraits virtual void computeLayerCells(const int iteration); virtual void findCellsNeighbours(const int iteration); virtual void findRoads(const int iteration); - virtual bool supportsExtendTracks() const noexcept { return true; } - virtual void extendTracks(const int iteration); template void processNeighbours(int iteration, int defaultCellTopologyId, int iLevel, const bounded_vector& currentCellSeed, const bounded_vector& currentCellId, const bounded_vector& currentCellTopologyId, bounded_vector& updatedCellSeed, bounded_vector& updatedCellId, bounded_vector& updatedCellTopologyId); - void acceptTracks(int iteration, bounded_vector& tracks, bounded_vector>& firstClusters); + void acceptTracks(int iteration, bounded_vector& tracks, bounded_vector>& firstClusters, size_t& nExtendedTracks, size_t& nExtendedClusters); void markTracks(int iteration); void updateTrackingParameters(const std::vector& trkPars) @@ -93,54 +92,19 @@ class TrackerTraits std::shared_ptr mTaskArena; protected: - using TrackExtensionCandidateN = TrackExtensionCandidate; - struct TrackExtensionCandidates { - TrackExtensionCandidates() = default; - explicit TrackExtensionCandidates(size_t nTracks) - : candidates(nTracks * MaxTrackExtensionCandidatesPerTrack), counts(nTracks, 0) - { - } - - int add(int trackIndex, const TrackExtensionCandidateN& candidate) - { - auto& count = counts[trackIndex]; - if (count >= MaxTrackExtensionCandidatesPerTrack) { - return -1; - } - const int flatIndex = static_cast(getFlatTrackExtensionCandidateIndex(trackIndex, count)); - candidates[flatIndex] = candidate; - ++count; - return flatIndex; - } - - void pop_back(int trackIndex) - { - --counts[trackIndex]; - } - - bool empty(int trackIndex) const { return counts[trackIndex] == 0; } - int size(int trackIndex) const { return counts[trackIndex]; } - TrackExtensionCandidateN* begin(int trackIndex) { return candidates.data() + getFlatTrackExtensionCandidateIndex(trackIndex, 0); } - TrackExtensionCandidateN* end(int trackIndex) { return begin(trackIndex) + counts[trackIndex]; } - TrackExtensionCandidateN& get(int trackIndex, int candidateIndex) { return candidates[getFlatTrackExtensionCandidateIndex(trackIndex, candidateIndex)]; } - const TrackExtensionCandidateN& get(int trackIndex, int candidateIndex) const { return candidates[getFlatTrackExtensionCandidateIndex(trackIndex, candidateIndex)]; } - TrackExtensionCandidateN& getFlat(int flatIndex) { return candidates[flatIndex]; } - - std::vector candidates; - std::vector counts; - }; - struct TrackFollowerScratch { std::vector> activeHypotheses; std::vector> nextHypotheses; }; + bool finaliseTrackSeed(const TrackSeedN& seed, + TrackITSExt& track, + const int iteration, + const TrackingFrameInfo* const* tfInfos, + const Cluster* const* unsortedClusters, + const o2::base::Propagator* propagator); bool trackFollowing(TrackITSExt* track, bool outward, const int iteration, TrackFollowerScratch& scratch); bool refitExtendedTrack(TrackITSExt& track, const int iteration); - void updateExtendedTrackTimeStamp(TrackITSExt& track, const int iteration); - virtual bool materializeTrackExtensionCandidate(TrackITSExt& track, const TrackExtensionCandidateN& candidate, const int iteration); - virtual void buildTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack); - void applyTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack); o2::gpu::GPUChainITS* mChain = nullptr; TimeFrame* mTimeFrame; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h index 054972faa8ed8..d80974e90a4ac 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h @@ -96,7 +96,7 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper void TimeFrame::wipe() { deepVectorClear(mTracks); - deepVectorClear(mFittedExtensionTracks); deepVectorClear(mTracklets); deepVectorClear(mCells); deepVectorClear(mCellsNeighbours); diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index 1a1b24cb8d78a..57c99f2557840 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -93,18 +93,6 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e logger(std::format(" - Neighbours finding: {} neighbours found in {:.2f} ms", nNeighbours, timeNeighbours)); logger(std::format(" - Track finding: {} tracks found in {:.2f} ms", nTracks + mTimeFrame->getNumberOfTracks(), timeRoads)); total += timeTracklets + timeCells + timeNeighbours + timeRoads; - if (mTraits->supportsExtendTracks() && (mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot])) { - const int nClustersBefore = mTimeFrame->getNumberOfUsedClusters(); - const int nTracksBefore = std::count_if(mTimeFrame->getTracks().begin(), mTimeFrame->getTracks().end(), [](const auto& track) { - return track.getPattern() & 0xff000000; - }); - const auto timeExtending = evaluateTask(&Tracker::extendTracks, StateNames[mCurStep = Extending], iteration, evalLog, iteration); - const int nTracksAfter = std::count_if(mTimeFrame->getTracks().begin(), mTimeFrame->getTracks().end(), [](const auto& track) { - return track.getPattern() & 0xff000000; - }); - total += timeExtending; - logger(std::format(" - Extending tracks: {} tracks using {} clusters in {:.2f} ms", nTracksAfter - nTracksBefore, mTimeFrame->getNumberOfUsedClusters() - nClustersBefore, timeExtending)); - } } } catch (const BoundedMemoryResource::MemoryLimitExceeded& err) { handleException(err); diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 39c834ca3ec55..7451fb3bff0a5 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -15,9 +15,10 @@ #include #include +#include #include +#include #include -#include #include #include #include @@ -663,6 +664,88 @@ void TrackerTraits::processNeighbours(int iteration, int defaultCellTop }); } +template +bool TrackerTraits::finaliseTrackSeed(const TrackSeedN& seed, + TrackITSExt& track, + const int iteration, + const TrackingFrameInfo* const* tfInfos, + const Cluster* const* unsortedClusters, + const o2::base::Propagator* propagator) +{ + if (!track::refitTrack(seed, + track, + mTrkParams[iteration].MaxChi2ClusterAttachment, + mTrkParams[iteration].MaxChi2NDF, + mBz, + tfInfos, + unsortedClusters, + mTrkParams[iteration].LayerxX0.data(), + mTrkParams[iteration].LayerRadii.data(), + mTrkParams[iteration].MinPt.data(), + propagator, + mTrkParams[iteration].CorrType, + mTrkParams[iteration].ReseedIfShorter, + mTrkParams[iteration].ShiftRefToCluster, + mTrkParams[iteration].RepeatRefitOut)) { + return false; + } + + const bool extendTop = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; + const bool extendBot = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; + if (!extendTop && !extendBot) { + return true; + } + + const auto backup = track; + auto best = track; + TrackFollowerScratch scratch; + const uint32_t lastLayer = static_cast(mTrkParams[iteration].NLayers - 1); + + auto finaliseExtensionCandidate = [&](TrackITSExt& candidate) { + const auto diff = (candidate.getPattern() & ~backup.getPattern()) & makeAddedClustersPatternMask(); + if (!diff || !refitExtendedTrack(candidate, iteration)) { + return; + } + applyExtendedClustersPattern(candidate, diff); + if (track::isBetter(candidate, best)) { + best = candidate; + } + }; + + std::optional topResult, botResult; + if (extendTop && backup.getLastClusterLayer() != lastLayer) { + auto candidate = backup; + if (trackFollowing(&candidate, true, iteration, scratch)) { + topResult = candidate; + finaliseExtensionCandidate(candidate); + } + } + if (extendBot && backup.getFirstClusterLayer() != 0) { + auto candidate = backup; + if (trackFollowing(&candidate, false, iteration, scratch)) { + botResult = candidate; + finaliseExtensionCandidate(candidate); + } + } + if (extendTop && extendBot) { + if (topResult && topResult->getFirstClusterLayer() != 0) { + auto candidate = *topResult; + if (trackFollowing(&candidate, false, iteration, scratch)) { + finaliseExtensionCandidate(candidate); + } + } + if (botResult && botResult->getLastClusterLayer() != lastLayer) { + auto candidate = *botResult; + if (trackFollowing(&candidate, true, iteration, scratch)) { + finaliseExtensionCandidate(candidate); + } + } + } + + track = best; + return true; +} + template void TrackerTraits::findRoads(const int iteration) { @@ -675,6 +758,7 @@ void TrackerTraits::findRoads(const int iteration) tfInfos[iLayer] = mTimeFrame->getTrackingFrameInfoOnLayer(iLayer).data(); unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } + size_t nExtendedTracks{0}, nExtendedClusters{0}; const auto topology = mTimeFrame->getTrackingTopologyView(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { @@ -723,65 +807,34 @@ void TrackerTraits::findRoads(const int iteration) bounded_vector tracks(mMemoryPool.get()); mTaskArena->execute([&] { - auto forSeed = [&](auto Tag, int iSeed, int offset = 0) { - TrackITSExt temporaryTrack; - bool refitSuccess = track::refitTrack(trackSeeds[iSeed], - temporaryTrack, - mTrkParams[iteration].MaxChi2ClusterAttachment, - mTrkParams[iteration].MaxChi2NDF, - mBz, - tfInfos, - unsortedClusters, - mTrkParams[iteration].LayerxX0.data(), - mTrkParams[iteration].LayerRadii.data(), - mTrkParams[iteration].MinPt.data(), - propagator, - mTrkParams[iteration].CorrType, - mTrkParams[iteration].ReseedIfShorter, - mTrkParams[iteration].ShiftRefToCluster, - mTrkParams[iteration].RepeatRefitOut); - - if (refitSuccess) { - if constexpr (decltype(Tag)::value == PassMode::OnePass::value) { - tracks.push_back(temporaryTrack); - } else if constexpr (decltype(Tag)::value == PassMode::TwoPassCount::value) { - // nothing to do - } else if constexpr (decltype(Tag)::value == PassMode::TwoPassInsert::value) { - tracks[offset] = temporaryTrack; - } else { - static_assert(false, "Unknown mode!"); - } - return 1; - } - return 0; - }; - const int nSeeds = static_cast(trackSeeds.size()); - if (mTaskArena->max_concurrency() <= 1) { - for (int iSeed{0}; iSeed < nSeeds; ++iSeed) { - forSeed(PassMode::OnePass{}, iSeed); - } - } else { - // The double-pass allows us to avoid sizeable memory spikes - bounded_vector perSeedCount(nSeeds + 1, 0, mMemoryPool.get()); - tbb::parallel_for(0, nSeeds, [&](const int iSeed) { - perSeedCount[iSeed] = forSeed(PassMode::TwoPassCount{}, iSeed); - }); - - std::exclusive_scan(perSeedCount.begin(), perSeedCount.end(), perSeedCount.begin(), 0); - auto totalTracks{perSeedCount.back()}; - if (totalTracks == 0) { - return; - } - tracks.resize(totalTracks); - - tbb::parallel_for(0, nSeeds, [&](const int iSeed) { - if (perSeedCount[iSeed] == perSeedCount[iSeed + 1]) { - return; + const int nWorkers = std::min(static_cast(mTaskArena->max_concurrency()), nSeeds); + const int chunkSize = std::min(nSeeds, std::clamp(nSeeds / (16 * nWorkers), 256, 4096)); + std::atomic nextSeed{0}; + std::mutex tracksMutex; + tbb::parallel_for(0, nWorkers, [&](const int) { + bounded_vector localTracks(mMemoryPool.get()); + localTracks.reserve(chunkSize); + while (true) { + const int firstSeed = nextSeed.fetch_add(chunkSize, std::memory_order_relaxed); + if (firstSeed >= nSeeds) { + break; } - forSeed(PassMode::TwoPassInsert{}, iSeed, perSeedCount[iSeed]); - }); - } + const int lastSeed = std::min(firstSeed + chunkSize, nSeeds); + for (int iSeed{firstSeed}; iSeed < lastSeed; ++iSeed) { + TrackITSExt temporaryTrack; + if (finaliseTrackSeed(trackSeeds[iSeed], temporaryTrack, iteration, tfInfos, unsortedClusters, propagator)) { + localTracks.push_back(temporaryTrack); + } + } + if (!localTracks.empty()) { + std::lock_guard lock{tracksMutex}; + tracks.insert(tracks.end(), std::make_move_iterator(localTracks.begin()), std::make_move_iterator(localTracks.end())); + localTracks.clear(); + } + } + deepVectorClear(localTracks); + }); deepVectorClear(trackSeeds); }); @@ -790,13 +843,20 @@ void TrackerTraits::findRoads(const int iteration) return track::isBetter(a, b); }); - acceptTracks(iteration, tracks, firstClusters); + acceptTracks(iteration, tracks, firstClusters, nExtendedTracks, nExtendedClusters); + } + if (mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]) { + LOGP(info, "Integrated track extension accepted {} tracks using {} clusters in iteration {}", nExtendedTracks, nExtendedClusters, iteration); } markTracks(iteration); } template -void TrackerTraits::acceptTracks(int iteration, bounded_vector& tracks, bounded_vector>& firstClusters) +void TrackerTraits::acceptTracks(int iteration, + bounded_vector& tracks, + bounded_vector>& firstClusters, + size_t& nExtendedTracks, + size_t& nExtendedClusters) { auto& trks = mTimeFrame->getTracks(); trks.reserve(trks.size() + tracks.size()); @@ -857,8 +917,14 @@ void TrackerTraits::acceptTracks(int iteration, bounded_vector smallestROFHalf) { track.getTimeStamp().setTimeStampError(smallestROFHalf); } - track.setUserField(0); - track.getParamOut().setUserField(0); + const auto diff = getAddedClustersPattern(track); + if (diff) { + ++nExtendedTracks; + for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { + nExtendedClusters += static_cast(diff & (0x1u << iLayer)); + } + } + clearAddedClustersPattern(track); trks.emplace_back(track); if (mTrkParams[iteration].AllowSharingFirstCluster) { @@ -913,17 +979,6 @@ void TrackerTraits::markTracks(int iteration) } } -template -void TrackerTraits::extendTracks(const int iteration) -{ - const auto nTracks = mTimeFrame->getTracks().size(); - TrackExtensionCandidates candidatesPerTrack(nTracks); - mTimeFrame->mFittedExtensionTracks.clear(); - buildTrackExtensionCandidates(iteration, candidatesPerTrack); - applyTrackExtensionCandidates(iteration, candidatesPerTrack); - mTimeFrame->mFittedExtensionTracks.clear(); -} - template bool TrackerTraits::refitExtendedTrack(TrackITSExt& track, const int iteration) { @@ -977,247 +1032,6 @@ bool TrackerTraits::refitExtendedTrack(TrackITSExt& track, const int it return fitSuccess; } -template -void TrackerTraits::updateExtendedTrackTimeStamp(TrackITSExt& track, const int iteration) -{ - bool firstCluster{true}, nominalCompatible{true}; - TimeEstBC nominalTS, expandedTS; - for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { - const int cluster = track.getClusterIndex(iLayer); - if (cluster == constants::UnusedIndex) { - continue; - } - const int rof = mTimeFrame->getClusterROF(iLayer, cluster); - const auto nominalROFTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(rof); - const auto expandedROFTS = mTimeFrame->getROFOverlapTableView().getLayer(iLayer).getROFTimeBounds(rof, true); - if (firstCluster) { - firstCluster = false; - nominalTS = nominalROFTS; - expandedTS = expandedROFTS; - continue; - } - if (nominalCompatible) { - if (nominalTS.isCompatible(nominalROFTS)) { - nominalTS += nominalROFTS; - } else { - nominalCompatible = false; - } - } - if (!expandedTS.isCompatible(expandedROFTS)) { - LOGP(fatal, "Clusters of an accepted track have non-overlapping expanded ROF time windows: {}+/-{} vs {}+/-{}", expandedROFTS.getTimeStamp(), expandedROFTS.getTimeStampError(), expandedTS.getTimeStamp(), expandedTS.getTimeStampError()); - } - expandedTS += expandedROFTS; - } - track.getTimeStamp() = (nominalCompatible ? nominalTS : expandedTS).makeSymmetrical(); -} - -template -bool TrackerTraits::materializeTrackExtensionCandidate(TrackITSExt& track, const TrackExtensionCandidateN& candidate, const int /*iteration*/) -{ - if (candidate.resultIndex < 0 || candidate.resultIndex >= static_cast(mTimeFrame->mFittedExtensionTracks.size())) { - return false; - } - track = mTimeFrame->mFittedExtensionTracks[candidate.resultIndex]; - return true; -} - -template -void TrackerTraits::buildTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack) -{ - struct ThreadExtensionResults { - std::vector tracks; - std::vector candidateIndicesToPatch; - }; - tbb::enumerable_thread_specific fittedTracks; - - auto prepareCandidate = [&](int trackIndex, const TrackITSExt& backup, TrackITSExt& candidate) { - if (!refitExtendedTrack(candidate, iteration)) { - return; - } - updateExtendedTrackTimeStamp(candidate, iteration); - const auto diff = (candidate.getPattern() & ~backup.getPattern()) & makeAddedClustersPatternMask(); - if (!diff) { - return; - } - applyExtendedClustersPattern(candidate, diff); - - TrackExtensionCandidateN extension; - extension.trackIndex = trackIndex; - for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { - if (backup.getClusterIndex(iLayer) == constants::UnusedIndex && candidate.getClusterIndex(iLayer) != constants::UnusedIndex) { - extension.addedClusters[iLayer] = candidate.getClusterIndex(iLayer); - ++extension.nAddedClusters; - } - } - if (!extension.nAddedClusters) { - return; - } - extension.chi2 = candidate.getChi2(); - const int candidateIndex = candidatesPerTrack.add(trackIndex, extension); - if (candidateIndex < 0) { - return; - } - auto& storedExtension = candidatesPerTrack.getFlat(candidateIndex); - auto& localFittedTracks = fittedTracks.local(); - storedExtension.resultIndex = static_cast(localFittedTracks.tracks.size()); - localFittedTracks.tracks.push_back(candidate); - localFittedTracks.candidateIndicesToPatch.push_back(candidateIndex); - }; - - const bool extendTop = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; - const bool extendBot = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; - auto& tracks = mTimeFrame->getTracks(); - tbb::enumerable_thread_specific trackFollowerScratch; - - const uint32_t lastLayer = static_cast(mTrkParams[iteration].NLayers - 1); - - auto buildCandidates = [&](int iTrack) { - const auto& backup = tracks[iTrack]; - auto& scratch = trackFollowerScratch.local(); - - std::optional topResult, botResult; - - if (extendTop && backup.getLastClusterLayer() != lastLayer) { - auto candidate{backup}; - if (trackFollowing(&candidate, true, iteration, scratch)) { - topResult = candidate; - prepareCandidate(iTrack, backup, candidate); - } - } - if (extendBot && backup.getFirstClusterLayer() != 0) { - auto candidate{backup}; - if (trackFollowing(&candidate, false, iteration, scratch)) { - botResult = candidate; - prepareCandidate(iTrack, backup, candidate); - } - } - if (extendTop && extendBot) { - if (topResult && topResult->getFirstClusterLayer() != 0) { - auto candidate = *topResult; - if (trackFollowing(&candidate, false, iteration, scratch)) { - prepareCandidate(iTrack, backup, candidate); - } - } - if (botResult && botResult->getLastClusterLayer() != lastLayer) { - auto candidate = *botResult; - if (trackFollowing(&candidate, true, iteration, scratch)) { - prepareCandidate(iTrack, backup, candidate); - } - } - } - }; - - if (mTaskArena->max_concurrency() <= 1) { - for (int iTrack{0}; iTrack < static_cast(tracks.size()); ++iTrack) { - buildCandidates(iTrack); - } - } else { - mTaskArena->execute([&] { - tbb::parallel_for(0, static_cast(tracks.size()), buildCandidates); - }); - } - - size_t nFittedExtensionTracks{0}; - for (auto& localFittedTracks : fittedTracks) { - nFittedExtensionTracks += localFittedTracks.tracks.size(); - } - mTimeFrame->mFittedExtensionTracks.reserve(nFittedExtensionTracks); - - int resultOffset{0}; - for (auto& localFittedTracks : fittedTracks) { - for (auto candidateIndex : localFittedTracks.candidateIndicesToPatch) { - candidatesPerTrack.getFlat(candidateIndex).resultIndex += resultOffset; - } - mTimeFrame->mFittedExtensionTracks.insert(mTimeFrame->mFittedExtensionTracks.end(), localFittedTracks.tracks.begin(), localFittedTracks.tracks.end()); - resultOffset += static_cast(localFittedTracks.tracks.size()); - } -} - -template -void TrackerTraits::applyTrackExtensionCandidates(const int iteration, TrackExtensionCandidates& candidatesPerTrack) -{ - auto& tracks = mTimeFrame->getTracks(); - - for (int iTrack{0}; iTrack < static_cast(tracks.size()); ++iTrack) { - std::stable_sort(candidatesPerTrack.begin(iTrack), candidatesPerTrack.end(iTrack), isBetterTrackExtensionCandidate); - while (!candidatesPerTrack.empty(iTrack) && (candidatesPerTrack.get(iTrack, candidatesPerTrack.size(iTrack) - 1).nAddedClusters <= 0)) { - candidatesPerTrack.pop_back(iTrack); - } - } - - std::array, NLayers> claimedClusters; - for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { - claimedClusters[iLayer].resize(mTimeFrame->getClusters()[iLayer].size(), 0); - } - - struct Entry { - int track; - int idx; - }; - auto cmp = [&](const Entry& a, const Entry& b) { - const auto& ca = candidatesPerTrack.get(a.track, a.idx); - const auto& cb = candidatesPerTrack.get(b.track, b.idx); - if (isBetterTrackExtensionCandidate(cb, ca)) { - return true; - } - if (isBetterTrackExtensionCandidate(ca, cb)) { - return false; - } - if (a.track != b.track) { - return a.track > b.track; - } - return a.idx > b.idx; - }; - std::priority_queue, decltype(cmp)> pq(cmp); - for (int iTrack{0}; iTrack < static_cast(tracks.size()); ++iTrack) { - if (!candidatesPerTrack.empty(iTrack)) { - pq.push({iTrack, 0}); - } - } - - auto tryNext = [&](int trackIndex, int idx) { - if (idx + 1 < candidatesPerTrack.size(trackIndex)) { - pq.push({trackIndex, idx + 1}); - } - }; - - while (!pq.empty()) { - const Entry e = pq.top(); - pq.pop(); - const auto& candidate = candidatesPerTrack.get(e.track, e.idx); - - bool hasContention{false}; - for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { - const int cluster = candidate.addedClusters[iLayer]; - if (cluster == constants::UnusedIndex) { - continue; - } - if (cluster >= static_cast(claimedClusters[iLayer].size()) || claimedClusters[iLayer][cluster]) { - hasContention = true; - break; - } - } - if (hasContention) { - tryNext(e.track, e.idx); - continue; - } - auto extendedTrack = tracks[e.track]; - if (!materializeTrackExtensionCandidate(extendedTrack, candidate, iteration)) { - tryNext(e.track, e.idx); - continue; - } - tracks[e.track] = extendedTrack; - for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { - const int cluster = candidate.addedClusters[iLayer]; - if (cluster == constants::UnusedIndex) { - continue; - } - claimedClusters[iLayer][cluster] = 1; - mTimeFrame->markUsedCluster(iLayer, cluster); - } - } -} - template bool TrackerTraits::trackFollowing(TrackITSExt* track, bool outward, const int iteration, TrackFollowerScratch& scratch) { From 2a2b4b9c1376521c5cf04a1a2b6120fe868f3838 Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Mon, 1 Jun 2026 22:20:47 +0200 Subject: [PATCH 3/4] ITS: implement code review --- .../ITS/include/DataFormatsITS/TrackITS.h | 42 +- .../GPU/ITStrackingGPU/TimeFrameGPU.h | 19 +- .../GPU/ITStrackingGPU/TrackingKernels.h | 54 +- .../ITS/tracking/GPU/cuda/TimeFrameGPU.cu | 62 -- .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 13 +- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 575 ++---------------- .../include/ITStracking/IndexTableUtils.h | 25 +- .../include/ITStracking/ROFLookupTables.h | 34 +- .../tracking/include/ITStracking/TimeFrame.h | 20 + .../ITStracking/TrackExtensionCandidate.h | 134 ---- .../ITStracking/TrackExtensionHypothesis.h | 56 ++ .../include/ITStracking/TrackFollower.h | 120 +--- .../include/ITStracking/TrackHelpers.h | 202 +++--- .../include/ITStracking/TrackITSInternal.h | 113 ++++ .../include/ITStracking/TrackerTraits.h | 21 +- .../include/ITStracking/TrackingConfigParam.h | 3 +- .../ITSMFT/ITS/tracking/src/Configuration.cxx | 13 +- .../ITSMFT/ITS/tracking/src/TimeFrame.cxx | 1 + Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx | 4 + .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 126 ++-- 20 files changed, 550 insertions(+), 1087 deletions(-) delete mode 100644 Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h create mode 100644 Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionHypothesis.h create mode 100644 Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackITSInternal.h diff --git a/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h b/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h index 20fb7c63ebacd..9b63509cc9424 100644 --- a/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h +++ b/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h @@ -35,6 +35,11 @@ namespace its class TrackITS : public o2::track::TrackParCov { + public: + static constexpr unsigned int ExtendedPatternShift = 24; + static constexpr int MaxLayersInTrackPattern = 8; + + private: enum UserBits { kSharedClusters = 1 << 28 }; @@ -106,8 +111,39 @@ class TrackITS : public o2::track::TrackParCov GPUhdi() uint32_t getPattern() const { return mPattern; } bool hasHitOnLayer(uint32_t i) const { return mPattern & (0x1 << i); } bool isFakeOnLayer(uint32_t i) const { return !(mPattern & (0x1 << (16 + i))); } - bool isExtendedOnLayer(uint32_t i) const { return (mPattern & (0x1 << (24 + i))); } // only correct if getNClusters <= 8 on layers <= 8 - uint32_t getLastClusterLayer() const + bool isExtendedOnLayer(uint32_t i) const { return (mPattern & (0x1 << (ExtendedPatternShift + i))); } // only correct if getNClusters <= 8 on layers <= 8 + template + GPUhdi() static constexpr uint32_t getLayerPatternMask() + { + return (NLayers >= 32) ? 0xffffffffu : ((1u << NLayers) - 1u); + } + template + GPUhdi() void setExtendedLayerPattern(uint32_t pattern) + { + pattern &= getLayerPatternMask(); + setUserField(static_cast(pattern)); + if constexpr (NLayers <= MaxLayersInTrackPattern) { + setPattern(getPattern() | (pattern << ExtendedPatternShift)); + } + } + template + GPUhdi() uint32_t getExtendedLayerPattern() const + { + const auto mask = getLayerPatternMask(); + if constexpr (NLayers <= MaxLayersInTrackPattern) { + const auto pattern = (getPattern() >> ExtendedPatternShift) & mask; + if (pattern) { + return pattern; + } + } + return getUserField() & mask; + } + GPUhdi() void clearExtendedLayerPattern() + { + setUserField(0); + getParamOut().setUserField(0); + } + GPUhdi() uint32_t getLastClusterLayer() const { uint32_t r{0}, v{mPattern & ((1 << 16) - 1)}; while (v >>= 1) { @@ -115,7 +151,7 @@ class TrackITS : public o2::track::TrackParCov } return r; } - uint32_t getFirstClusterLayer() const + GPUhdi() uint32_t getFirstClusterLayer() const { int s{0}; while (!(mPattern & (1 << s))) { diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index 7223968c8cbf9..3f574ff51a8c3 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -19,7 +19,7 @@ #include "ITStracking/BoundedAllocator.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/Configuration.h" -#include "ITStracking/TrackExtensionCandidate.h" +#include "ITStracking/TrackExtensionHypothesis.h" #include "ITStrackingGPU/Utils.h" namespace o2::its::gpu @@ -91,13 +91,9 @@ class TimeFrameGPU : public TimeFrame void createNeighboursDevice(const unsigned int layer); void createNeighboursLUTDevice(const int, const unsigned int); void createTrackITSExtDevice(const size_t); - void loadTrackExtensionStartTracksDevice(); - void createTrackExtensionCandidatesDevice(const size_t); void createTrackExtensionScratchDevice(const int nThreads, const int beamWidth); - void createTrackExtensionResultsDevice(const size_t); void downloadTrackITSExtDevice(); void downloadCellsNeighboursDevice(std::vector>&, const int); - void downloadTrackExtensionResultsDevice(); void downloadNeighboursLUTDevice(bounded_vector&, const int); void downloadCellsDevice(); void downloadCellsLUTDevice(); @@ -124,7 +120,6 @@ class TimeFrameGPU : public TimeFrame const auto getDeviceTrackingTopologyView() const { return mDeviceTrackingTopologyView; } int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; } auto& getTrackITSExt() { return mTrackITSExt; } - auto& getTrackExtensionResults() { return mTrackExtensionResults; } Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; } int* getDeviceROFramesPV() { return mROFramesPVDevice; } unsigned char* getDeviceUsedClusters(const int); @@ -132,12 +127,8 @@ class TimeFrameGPU : public TimeFrame // Hybrid TrackITSExt* getDeviceTrackITSExt() { return mTrackITSExtDevice; } - TrackITSExt* getDeviceTrackExtensionStartTracks() { return mTrackExtensionStartTracksDevice; } - TrackExtensionCandidate* getDeviceTrackExtensionCandidates() { return mTrackExtensionCandidatesDevice; } - int* getDeviceTrackExtensionCandidateOffsets() { return mTrackExtensionCandidateOffsetsDevice; } TrackExtensionHypothesis* getDeviceActiveTrackExtensionHypotheses() { return mActiveTrackExtensionHypothesesDevice; } TrackExtensionHypothesis* getDeviceNextTrackExtensionHypotheses() { return mNextTrackExtensionHypothesesDevice; } - TrackExtensionResult* getDeviceTrackExtensionResults() { return mTrackExtensionResultsDevice; } int* getDeviceNeighboursLUT(const int layer) { return mNeighboursLUTDevice[layer]; } gsl::span getDeviceNeighboursLUTs() { return mNeighboursLUTDevice; } CellNeighbour** getDeviceArrayNeighbours() { return mNeighboursDeviceArray; } @@ -235,13 +226,8 @@ class TimeFrameGPU : public TimeFrame float** mCellSeedsChi2DeviceArray; TrackITSExt* mTrackITSExtDevice; - TrackITSExt* mTrackExtensionStartTracksDevice{nullptr}; - TrackExtensionCandidate* mTrackExtensionCandidatesDevice{nullptr}; - int* mTrackExtensionCandidateOffsetsDevice{nullptr}; TrackExtensionHypothesis* mActiveTrackExtensionHypothesesDevice{nullptr}; TrackExtensionHypothesis* mNextTrackExtensionHypothesesDevice{nullptr}; - TrackExtensionResult* mTrackExtensionResultsDevice{nullptr}; - unsigned int mNTrackExtensionResults{0}; std::array mNeighboursDevice{}; CellNeighbour** mNeighboursDeviceArray{nullptr}; std::array mTrackingFrameInfoDevice; @@ -258,9 +244,6 @@ class TimeFrameGPU : public TimeFrame // Temporary buffer for storing output tracks from GPU tracking bounded_vector mTrackITSExt; - bounded_vector mTrackExtensionStartTracks; - // Temporary buffer for fitted track extension proposals from GPU tracking - bounded_vector> mTrackExtensionResults; }; template diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index ff541e0e5a839..4b3b36f513574 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -19,7 +19,7 @@ #include "ITStracking/BoundedAllocator.h" #include "ITStracking/ROFLookupTables.h" #include "ITStracking/TrackingTopology.h" -#include "ITStracking/TrackExtensionCandidate.h" +#include "ITStracking/TrackExtensionHypothesis.h" #include "ITStrackingGPU/Utils.h" #include "DetectorsBase/Propagator.h" @@ -37,58 +37,6 @@ class Cluster; class TrackITSExt; class ExternalAllocator; -inline constexpr int kTrackExtensionLaunchBlocks = 60; -inline constexpr int kTrackExtensionLaunchThreadsPerBlock = 256; -inline constexpr int kTrackExtensionLaunchThreads = kTrackExtensionLaunchBlocks * kTrackExtensionLaunchThreadsPerBlock; - -template -void computeTrackExtensionCandidatesHandler(const TrackITSExt* tracks, - const IndexTableUtils* utils, - const typename ROFMaskTable::View& rofMask, - const typename ROFOverlapTable::View& rofOverlaps, - const Cluster** clusters, - const unsigned char** usedClusters, - const int** clustersIndexTables, - const int** ROFClusters, - const TrackingFrameInfo** trackingFrameInfo, - TrackExtensionCandidate* candidates, - int* candidateOffsets, - TrackExtensionHypothesis* activeHypotheses, - TrackExtensionHypothesis* nextHypotheses, - const std::array layerRadii, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const int phiBins, - const int beamWidth, - const bool extendTop, - const bool extendBot, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const float nSigmaCutPhi, - const float nSigmaCutZ, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - gpu::Stream& stream); - -template -void computeTrackExtensionResultsHandler(const TrackITSExt* tracks, - const TrackExtensionCandidate* candidates, - const int* candidateOffsets, - TrackExtensionResult* results, - const TrackingFrameInfo** trackingFrameInfo, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster, - gpu::Stream& stream); - template void countTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index af6a86665de96..432ee62ab7b6e 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -582,35 +582,6 @@ void TimeFrameGPU::createTrackITSExtDevice(const size_t nSeeds) GPUChkErrS(cudaMemset(mTrackITSExtDevice, 0, mNTracks * sizeof(o2::its::TrackITSExt))); } -template -void TimeFrameGPU::loadTrackExtensionStartTracksDevice() -{ - GPUTimer timer("loading track extension start tracks"); - GPULog("gpu-transfer: loading {} track extension start tracks, for {:.2f} MB.", this->mTracks.size(), this->mTracks.size() * sizeof(o2::its::TrackITSExt) / constants::MB); - mTrackExtensionStartTracksDevice = nullptr; - mTrackExtensionStartTracks = bounded_vector(this->mTracks.begin(), this->mTracks.end(), this->getMemoryPool().get()); - if (this->mTracks.empty()) { - return; - } - allocMem(reinterpret_cast(&mTrackExtensionStartTracksDevice), mTrackExtensionStartTracks.size() * sizeof(o2::its::TrackITSExt), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); - GPUChkErrS(cudaMemcpy(mTrackExtensionStartTracksDevice, mTrackExtensionStartTracks.data(), mTrackExtensionStartTracks.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyHostToDevice)); -} - -template -void TimeFrameGPU::createTrackExtensionCandidatesDevice(const size_t nTracks) -{ - GPUTimer timer("reserving track extension candidates"); - const size_t nCandidates = nTracks * MaxTrackExtensionCandidatesPerTrack; - GPULog("gpu-allocation: reserving {} track extension candidates, for {:.2f} MB.", nCandidates, nCandidates * sizeof(o2::its::TrackExtensionCandidate) / constants::MB); - mTrackExtensionCandidatesDevice = nullptr; - mTrackExtensionCandidateOffsetsDevice = nullptr; - if (nCandidates == 0) { - return; - } - allocMem(reinterpret_cast(&mTrackExtensionCandidatesDevice), nCandidates * sizeof(o2::its::TrackExtensionCandidate), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); - allocMem(reinterpret_cast(&mTrackExtensionCandidateOffsetsDevice), (nTracks + 1) * sizeof(int), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); -} - template void TimeFrameGPU::createTrackExtensionScratchDevice(const int nThreads, const int beamWidth) { @@ -626,28 +597,6 @@ void TimeFrameGPU::createTrackExtensionScratchDevice(const int nThreads allocMem(reinterpret_cast(&mNextTrackExtensionHypothesesDevice), nHypotheses * sizeof(o2::its::TrackExtensionHypothesis), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); } -template -void TimeFrameGPU::createTrackExtensionResultsDevice(const size_t nTracks) -{ - GPUTimer timer("reserving fitted track extension results"); - mNTrackExtensionResults = 0; - if (nTracks == 0 || mTrackExtensionCandidateOffsetsDevice == nullptr) { - mTrackExtensionResults = bounded_vector>(0, {}, this->getMemoryPool().get()); - mTrackExtensionResultsDevice = nullptr; - return; - } - int nResults{0}; - GPUChkErrS(cudaMemcpy(&nResults, mTrackExtensionCandidateOffsetsDevice + nTracks, sizeof(int), cudaMemcpyDeviceToHost)); - mNTrackExtensionResults = nResults; - GPULog("gpu-allocation: reserving {} fitted track extension results, for {:.2f} MB.", mNTrackExtensionResults, mNTrackExtensionResults * sizeof(o2::its::TrackExtensionResult) / constants::MB); - mTrackExtensionResults = bounded_vector>(mNTrackExtensionResults, {}, this->getMemoryPool().get()); - mTrackExtensionResultsDevice = nullptr; - if (mTrackExtensionResults.empty()) { - return; - } - allocMem(reinterpret_cast(&mTrackExtensionResultsDevice), mNTrackExtensionResults * sizeof(o2::its::TrackExtensionResult), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); -} - template void TimeFrameGPU::downloadCellsDevice() { @@ -694,17 +643,6 @@ void TimeFrameGPU::downloadTrackITSExtDevice() GPUChkErrS(cudaMemcpy(mTrackITSExt.data(), mTrackITSExtDevice, mTrackITSExt.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyDeviceToHost)); } -template -void TimeFrameGPU::downloadTrackExtensionResultsDevice() -{ - GPUTimer timer("downloading fitted track extension results"); - GPULog("gpu-transfer: downloading {} fitted track extension results, for {:.2f} MB.", mTrackExtensionResults.size(), mTrackExtensionResults.size() * sizeof(o2::its::TrackExtensionResult) / constants::MB); - if (mTrackExtensionResults.empty()) { - return; - } - GPUChkErrS(cudaMemcpy(mTrackExtensionResults.data(), mTrackExtensionResultsDevice, mTrackExtensionResults.size() * sizeof(o2::its::TrackExtensionResult), cudaMemcpyDeviceToHost)); -} - template void TimeFrameGPU::unregisterHostMemory(const int maxLayers) { diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 43c45649b656a..4cacf94684104 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -21,6 +21,10 @@ namespace o2::its { +namespace +{ +constexpr int trackExtensionLaunchThreads = 60 * 256; +} template void TrackerTraitsGPU::initialiseTimeFrame(const int iteration) @@ -309,8 +313,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; - size_t nExtendedTracks{0}; - size_t nExtendedClusters{0}; 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) { @@ -369,7 +371,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) mTimeFrameGPU->getFrameworkAllocator()); mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size()); if (extendTracks) { - mTimeFrameGPU->createTrackExtensionScratchDevice(kTrackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerBeamWidth); + mTimeFrameGPU->createTrackExtensionScratchDevice(trackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerBeamWidth); } computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), @@ -409,12 +411,9 @@ void TrackerTraitsGPU::findRoads(const int iteration) mTimeFrameGPU->downloadTrackITSExtDevice(); auto& tracks = mTimeFrameGPU->getTrackITSExt(); - this->acceptTracks(iteration, tracks, firstClusters, nExtendedTracks, nExtendedClusters); + this->acceptTracks(iteration, tracks, firstClusters); mTimeFrameGPU->loadUsedClustersDevice(); } - if (extendTracks) { - LOGP(info, "Integrated track extension accepted {} tracks using {} clusters in iteration {}", nExtendedTracks, nExtendedClusters, iteration); - } this->markTracks(iteration); // wipe the artefact memory mTimeFrameGPU->popMemoryStack(iteration); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 55a0bc4d069e0..69a2ff5be56d6 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -111,304 +111,29 @@ struct compare_track_chi2 { }; template -GPUdi() void writeTrackExtensionCandidate(const int trackIndex, - const TrackITSExt& original, - const TrackITSExt& updated, - TrackExtensionCandidate* candidates, - int& slot) -{ - if (slot >= MaxTrackExtensionCandidatesPerTrack) { - return; - } - auto& candidate = candidates[getFlatTrackExtensionCandidateIndex(trackIndex, slot)]; - candidate.reset(); - candidate.trackIndex = trackIndex; - for (int iLayer{0}; iLayer < NLayers; ++iLayer) { - if (original.getClusterIndex(iLayer) == constants::UnusedIndex && updated.getClusterIndex(iLayer) != constants::UnusedIndex) { - candidate.addedClusters[iLayer] = updated.getClusterIndex(iLayer); - ++candidate.nAddedClusters; - } - } - if (!candidate.nAddedClusters) { - candidate.reset(); - return; - } - candidate.chi2 = updated.getChi2(); - ++slot; -} - -template -GPUg() void __launch_bounds__(256, 1) computeTrackExtensionCandidatesKernel(const TrackITSExt* tracks, - const IndexTableUtils* utils, - const typename ROFMaskTable::View rofMask, - const typename ROFOverlapTable::View rofOverlaps, - const Cluster** clusters, - const unsigned char** usedClusters, - const int** clustersIndexTables, - const int** ROFClusters, - const TrackingFrameInfo** trackingFrameInfo, - TrackExtensionCandidate* candidates, - int* candidateOffsets, - TrackExtensionHypothesis* activeHypothesesScratch, - TrackExtensionHypothesis* nextHypothesesScratch, - const std::array layerRadii, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const int phiBins, - const int beamWidth, - const bool extendTop, - const bool extendBot, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const float nSigmaCutPhi, - const float nSigmaCutZ, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType) -{ - if (blockIdx.x == 0 && threadIdx.x == 0) { - candidateOffsets[nTracks] = 0; - } - const int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x; - auto* const threadActiveHypotheses = activeHypothesesScratch + (globalThreadId * beamWidth); - auto* const threadNextHypotheses = nextHypothesesScratch + (globalThreadId * beamWidth); - for (int iTrack = globalThreadId; iTrack < nTracks; iTrack += blockDim.x * gridDim.x) { - for (int iCandidate{0}; iCandidate < MaxTrackExtensionCandidatesPerTrack; ++iCandidate) { - candidates[getFlatTrackExtensionCandidateIndex(iTrack, iCandidate)].reset(); - } - const auto& track = tracks[iTrack]; - auto* activeHypotheses = threadActiveHypotheses; - auto* nextHypotheses = threadNextHypotheses; - int slot{0}; - if (extendTop && getTrackExtensionLastClusterLayer(track) != nLayers - 1) { - TrackITSExt topCandidate; - if (followTrackExtensionDirection(track, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, true, propagator, matCorrType, activeHypotheses, nextHypotheses, topCandidate)) { - writeTrackExtensionCandidate(iTrack, track, topCandidate, candidates, slot); - if (extendBot && getTrackExtensionFirstClusterLayer(topCandidate) != 0) { - TrackITSExt topBottomCandidate; - if (followTrackExtensionDirection(topCandidate, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, false, propagator, matCorrType, activeHypotheses, nextHypotheses, topBottomCandidate)) { - writeTrackExtensionCandidate(iTrack, track, topBottomCandidate, candidates, slot); - } - } - } - } - if (extendBot && getTrackExtensionFirstClusterLayer(track) != 0) { - TrackITSExt bottomCandidate; - if (followTrackExtensionDirection(track, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, false, propagator, matCorrType, activeHypotheses, nextHypotheses, bottomCandidate)) { - writeTrackExtensionCandidate(iTrack, track, bottomCandidate, candidates, slot); - if (extendTop && getTrackExtensionLastClusterLayer(bottomCandidate) != nLayers - 1) { - TrackITSExt bottomTopCandidate; - if (followTrackExtensionDirection(bottomCandidate, *utils, rofMask, rofOverlaps, clusters, usedClusters, clustersIndexTables, ROFClusters, trackingFrameInfo, layerRadii.data(), layerxX0.data(), nLayers, phiBins, beamWidth, bz, maxChi2ClusterAttachment, maxChi2NDF, nSigmaCutPhi, nSigmaCutZ, true, propagator, matCorrType, activeHypotheses, nextHypotheses, bottomTopCandidate)) { - writeTrackExtensionCandidate(iTrack, track, bottomTopCandidate, candidates, slot); - } - } - } - } - candidateOffsets[iTrack] = slot; - } -} - -template -GPUdi() bool fitTrackExtensionResult(const TrackITSExt& startTrack, - const TrackExtensionCandidate& candidate, - const TrackingFrameInfo* const* trackingFrameInfo, - const float* layerxX0, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster, - TrackITSExt& track) -{ - track = startTrack; - for (int iLayer{0}; iLayer < nLayers; ++iLayer) { - if (candidate.addedClusters[iLayer] != constants::UnusedIndex) { - track.setExternalClusterIndex(iLayer, candidate.addedClusters[iLayer], true); - } - } - - o2::track::TrackPar linRef{track}; - o2::its::track::resetTrackCovariance(track); - track.setChi2(0); - bool fitSuccess = o2::its::track::fitTrack(track, - 0, - nLayers, - 1, - maxChi2ClusterAttachment, - maxChi2NDF, - o2::constants::math::VeryBig, - 0, - bz, - trackingFrameInfo, - layerxX0, - propagator, - matCorrType, - &linRef, - shiftRefToCluster); - if (!fitSuccess) { - return false; - } - - track.getParamOut() = track.getParamIn(); - linRef = track.getParamOut(); - o2::its::track::resetTrackCovariance(track); - track.setChi2(0); - fitSuccess = o2::its::track::fitTrack(track, - nLayers - 1, - -1, - -1, - maxChi2ClusterAttachment, - maxChi2NDF, - 50.f, - 0, - bz, - trackingFrameInfo, - layerxX0, - propagator, - matCorrType, - &linRef, - shiftRefToCluster); - if (!fitSuccess) { - return false; - } - - uint32_t diff{0}; - for (int iLayer{0}; iLayer < nLayers; ++iLayer) { - if (candidate.addedClusters[iLayer] != constants::UnusedIndex) { - diff |= (0x1u << iLayer); - } - } - applyExtendedClustersPattern(track, diff); - return true; -} - -template -GPUdi() bool refitTrackExtensionResult(TrackITSExt& track, - const TrackingFrameInfo* const* trackingFrameInfo, - const float* layerxX0, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster) -{ - o2::track::TrackPar linRef{track}; - o2::its::track::resetTrackCovariance(track); - track.setChi2(0); - bool fitSuccess = o2::its::track::fitTrack(track, - 0, - nLayers, - 1, - maxChi2ClusterAttachment, - maxChi2NDF, - o2::constants::math::VeryBig, - 0, - bz, - trackingFrameInfo, - layerxX0, - propagator, - matCorrType, - &linRef, - shiftRefToCluster); - if (!fitSuccess) { - return false; - } - - track.getParamOut() = track.getParamIn(); - linRef = track.getParamOut(); - o2::its::track::resetTrackCovariance(track); - track.setChi2(0); - return o2::its::track::fitTrack(track, - nLayers - 1, - -1, - -1, - maxChi2ClusterAttachment, - maxChi2NDF, - 50.f, - 0, - bz, - trackingFrameInfo, - layerxX0, - propagator, - matCorrType, - &linRef, - shiftRefToCluster); -} - -template -GPUdi() void finaliseTrackExtensionCandidate(const uint32_t backupPattern, - TrackITSExt& candidate, - const TrackingFrameInfo* const* trackingFrameInfo, - const float* layerxX0, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster, - TrackITSExt& best) +GPUdi() void finaliseTrackExtensionTrial(const uint32_t backupPattern, + TrackITSInternal& trial, + const TrackingFrameInfo* const* trackingFrameInfo, + const float* layerxX0, + const int nLayers, + const float bz, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const bool shiftRefToCluster, + const bool repeatRefitOut, + TrackITSInternal& best, + uint32_t& bestDiff) { - const auto diff = (candidate.getPattern() & ~backupPattern) & makeAddedClustersPatternMask(); - if (!diff || !refitTrackExtensionResult(candidate, trackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster)) { + const auto diff = (trial.getPattern() & ~backupPattern) & TrackITS::getLayerPatternMask(); + if (!diff || + !o2::its::track::refitTrack(trial, trackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, repeatRefitOut)) { return; } - applyExtendedClustersPattern(candidate, diff); - if (o2::its::track::isBetter(candidate, best)) { - best = candidate; - } -} - -template -GPUg() void __launch_bounds__(256, 1) computeTrackExtensionResultsKernel(const TrackITSExt* tracks, - const TrackExtensionCandidate* candidates, - const int* candidateOffsets, - TrackExtensionResult* results, - const TrackingFrameInfo** trackingFrameInfo, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster) -{ - for (int iTrack = blockIdx.x * blockDim.x + threadIdx.x; iTrack < nTracks; iTrack += blockDim.x * gridDim.x) { - const int firstResult = candidateOffsets[iTrack]; - const int nResults = candidateOffsets[iTrack + 1] - firstResult; - const auto& startTrack = tracks[iTrack]; - for (int iCandidate{0}; iCandidate < nResults; ++iCandidate) { - const auto& candidate = candidates[getFlatTrackExtensionCandidateIndex(iTrack, iCandidate)]; - auto& result = results[firstResult + iCandidate]; - result.reset(); - if (!candidate.isValidForTrack(iTrack)) { - continue; - } - result.candidate = candidate; - if (!fitTrackExtensionResult(startTrack, - candidate, - trackingFrameInfo, - layerxX0.data(), - nLayers, - bz, - maxChi2ClusterAttachment, - maxChi2NDF, - propagator, - matCorrType, - shiftRefToCluster, - result.track)) { - result.reset(); - continue; - } - result.candidate.chi2 = result.track.getChi2(); - } + if (o2::its::track::isBetter(trial, best)) { + best = trial; + bestDiff = diff; } } @@ -432,7 +157,7 @@ GPUg() void __launch_bounds__(256, 1) countTrackSeedsKernel( const o2::base::PropagatorF::MatCorrType matCorrType) { for (int iCurrentTrackSeedIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackSeedIndex < nSeeds; iCurrentTrackSeedIndex += blockDim.x * gridDim.x) { - TrackITSExt temporaryTrack; + TrackITSInternal temporaryTrack; if (o2::its::track::refitTrack(trackSeeds[iCurrentTrackSeedIndex], temporaryTrack, maxChi2ClusterAttachment, @@ -493,7 +218,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( if (seedLUT[iCurrentTrackSeedIndex] == seedLUT[iCurrentTrackSeedIndex + 1]) { continue; } - TrackITSExt temporaryTrack; + TrackITSInternal temporaryTrack; bool refitSuccess = o2::its::track::refitTrack(trackSeeds[iCurrentTrackSeedIndex], temporaryTrack, maxChi2ClusterAttachment, @@ -517,15 +242,18 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( auto* nextHypotheses = nextHypothesesScratch + threadIndex * beamWidth; const auto backupPattern = temporaryTrack.getPattern(); auto best = temporaryTrack; - TrackITSExt topResult; - TrackITSExt botResult; + uint32_t bestDiff{0}; + TrackITSInternal topResult; + TrackITSInternal botResult; bool hasTopResult{false}; bool hasBotResult{false}; const uint32_t lastLayer = static_cast(nLayers - 1); - if (extendTop && getTrackExtensionLastClusterLayer(temporaryTrack) != lastLayer) { + if (extendTop && temporaryTrack.getLastClusterLayer() != lastLayer) { auto candidate = temporaryTrack; - if (followTrackExtensionDirection(temporaryTrack, + const auto startHypothesis = TrackExtensionHypothesis{temporaryTrack, true}; + TrackExtensionHypothesis bestHypothesis; + if (followTrackExtensionDirection(startHypothesis, *utils, rofMask, rofOverlaps, @@ -549,15 +277,18 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( matCorrType, activeHypotheses, nextHypotheses, - candidate)) { + bestHypothesis)) { + updateTrackFromExtensionHypothesis(bestHypothesis, true, nLayers, candidate); topResult = candidate; hasTopResult = true; - finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + finaliseTrackExtensionTrial(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, repeatRefitOut, best, bestDiff); } } - if (extendBot && getTrackExtensionFirstClusterLayer(temporaryTrack) != 0) { + if (extendBot && temporaryTrack.getFirstClusterLayer() != 0) { auto candidate = temporaryTrack; - if (followTrackExtensionDirection(temporaryTrack, + const auto startHypothesis = TrackExtensionHypothesis{temporaryTrack, false}; + TrackExtensionHypothesis bestHypothesis; + if (followTrackExtensionDirection(startHypothesis, *utils, rofMask, rofOverlaps, @@ -581,16 +312,19 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( matCorrType, activeHypotheses, nextHypotheses, - candidate)) { + bestHypothesis)) { + updateTrackFromExtensionHypothesis(bestHypothesis, false, nLayers, candidate); botResult = candidate; hasBotResult = true; - finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + finaliseTrackExtensionTrial(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, repeatRefitOut, best, bestDiff); } } if (extendTop && extendBot) { - if (hasTopResult && getTrackExtensionFirstClusterLayer(topResult) != 0) { + if (hasTopResult && topResult.getFirstClusterLayer() != 0) { auto candidate = topResult; - if (followTrackExtensionDirection(topResult, + const auto startHypothesis = TrackExtensionHypothesis{topResult, false}; + TrackExtensionHypothesis bestHypothesis; + if (followTrackExtensionDirection(startHypothesis, *utils, rofMask, rofOverlaps, @@ -614,13 +348,16 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( matCorrType, activeHypotheses, nextHypotheses, - candidate)) { - finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + bestHypothesis)) { + updateTrackFromExtensionHypothesis(bestHypothesis, false, nLayers, candidate); + finaliseTrackExtensionTrial(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, repeatRefitOut, best, bestDiff); } } - if (hasBotResult && getTrackExtensionLastClusterLayer(botResult) != lastLayer) { + if (hasBotResult && botResult.getLastClusterLayer() != lastLayer) { auto candidate = botResult; - if (followTrackExtensionDirection(botResult, + const auto startHypothesis = TrackExtensionHypothesis{botResult, true}; + TrackExtensionHypothesis bestHypothesis; + if (followTrackExtensionDirection(startHypothesis, *utils, rofMask, rofOverlaps, @@ -644,14 +381,20 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( matCorrType, activeHypotheses, nextHypotheses, - candidate)) { - finaliseTrackExtensionCandidate(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, best); + bestHypothesis)) { + updateTrackFromExtensionHypothesis(bestHypothesis, true, nLayers, candidate); + finaliseTrackExtensionTrial(backupPattern, candidate, foundTrackingFrameInfo, layerxX0, nLayers, bz, maxChi2ClusterAttachment, maxChi2NDF, propagator, matCorrType, shiftRefToCluster, repeatRefitOut, best, bestDiff); } } } temporaryTrack = best; + tracks[seedLUT[iCurrentTrackSeedIndex]] = makeTrackITSExt(temporaryTrack); + if (bestDiff) { + tracks[seedLUT[iCurrentTrackSeedIndex]].setExtendedLayerPattern(bestDiff); + } + continue; } - tracks[seedLUT[iCurrentTrackSeedIndex]] = temporaryTrack; + tracks[seedLUT[iCurrentTrackSeedIndex]] = makeTrackITSExt(temporaryTrack); } } } @@ -1078,114 +821,6 @@ GPUg() void __launch_bounds__(256, 1) processNeighboursKernel( } // namespace gpu -template -void computeTrackExtensionCandidatesHandler(const TrackITSExt* tracks, - const IndexTableUtils* utils, - const typename ROFMaskTable::View& rofMask, - const typename ROFOverlapTable::View& rofOverlaps, - const Cluster** clusters, - const unsigned char** usedClusters, - const int** clustersIndexTables, - const int** ROFClusters, - const TrackingFrameInfo** trackingFrameInfo, - TrackExtensionCandidate* candidates, - int* candidateOffsets, - TrackExtensionHypothesis* activeHypotheses, - TrackExtensionHypothesis* nextHypotheses, - const std::array layerRadii, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const int phiBins, - const int beamWidth, - const bool extendTop, - const bool extendBot, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const float nSigmaCutPhi, - const float nSigmaCutZ, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - gpu::Stream& stream) -{ - if (nTracks <= 0 || candidates == nullptr || candidateOffsets == nullptr || activeHypotheses == nullptr || nextHypotheses == nullptr) { - return; - } - gpu::computeTrackExtensionCandidatesKernel<<>>( - tracks, - utils, - rofMask, - rofOverlaps, - clusters, - usedClusters, - clustersIndexTables, - ROFClusters, - trackingFrameInfo, - candidates, - candidateOffsets, - activeHypotheses, - nextHypotheses, - layerRadii, - layerxX0, - nTracks, - nLayers, - phiBins, - beamWidth, - extendTop, - extendBot, - bz, - maxChi2ClusterAttachment, - maxChi2NDF, - nSigmaCutPhi, - nSigmaCutZ, - propagator, - matCorrType); - GPUChkErrS(cudaGetLastError()); - GPUChkErrS(cudaStreamSynchronize(stream.get())); - thrust::device_ptr offsets(candidateOffsets); - thrust::exclusive_scan(offsets, offsets + nTracks + 1, offsets); -} - -template -void computeTrackExtensionResultsHandler(const TrackITSExt* tracks, - const TrackExtensionCandidate* candidates, - const int* candidateOffsets, - TrackExtensionResult* results, - const TrackingFrameInfo** trackingFrameInfo, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster, - gpu::Stream& stream) -{ - if (nTracks <= 0 || tracks == nullptr || candidates == nullptr || candidateOffsets == nullptr || results == nullptr) { - return; - } - gpu::computeTrackExtensionResultsKernel<<>>( - tracks, - candidates, - candidateOffsets, - results, - trackingFrameInfo, - layerxX0, - nTracks, - nLayers, - bz, - maxChi2ClusterAttachment, - maxChi2NDF, - propagator, - matCorrType, - shiftRefToCluster); - GPUChkErrS(cudaGetLastError()); - GPUChkErrS(cudaStreamSynchronize(stream.get())); -} - template void countTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, @@ -1760,52 +1395,6 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, } /// Explicit instantiation of ITS2 handlers -template void computeTrackExtensionCandidatesHandler<7>(const TrackITSExt* tracks, - const IndexTableUtils<7>* utils, - const ROFMaskTable<7>::View& rofMask, - const ROFOverlapTable<7>::View& rofOverlaps, - const Cluster** clusters, - const unsigned char** usedClusters, - const int** clustersIndexTables, - const int** ROFClusters, - const TrackingFrameInfo** trackingFrameInfo, - TrackExtensionCandidate<7>* candidates, - int* candidateOffsets, - TrackExtensionHypothesis<7>* activeHypotheses, - TrackExtensionHypothesis<7>* nextHypotheses, - const std::array layerRadii, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const int phiBins, - const int beamWidth, - const bool extendTop, - const bool extendBot, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const float nSigmaCutPhi, - const float nSigmaCutZ, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - gpu::Stream& stream); - -template void computeTrackExtensionResultsHandler<7>(const TrackITSExt* tracks, - const TrackExtensionCandidate<7>* candidates, - const int* candidateOffsets, - TrackExtensionResult<7>* results, - const TrackingFrameInfo** trackingFrameInfo, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster, - gpu::Stream& stream); - template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const ROFMaskTable<7>::View& rofMask, const int transitionId, @@ -2006,52 +1595,6 @@ template void computeTrackSeedHandler(TrackSeed<7>* trackSeeds, /// Explicit instantiation of ALICE3 handlers #ifdef ENABLE_UPGRADES -template void computeTrackExtensionCandidatesHandler<11>(const TrackITSExt* tracks, - const IndexTableUtils<11>* utils, - const ROFMaskTable<11>::View& rofMask, - const ROFOverlapTable<11>::View& rofOverlaps, - const Cluster** clusters, - const unsigned char** usedClusters, - const int** clustersIndexTables, - const int** ROFClusters, - const TrackingFrameInfo** trackingFrameInfo, - TrackExtensionCandidate<11>* candidates, - int* candidateOffsets, - TrackExtensionHypothesis<11>* activeHypotheses, - TrackExtensionHypothesis<11>* nextHypotheses, - const std::array layerRadii, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const int phiBins, - const int beamWidth, - const bool extendTop, - const bool extendBot, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const float nSigmaCutPhi, - const float nSigmaCutZ, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - gpu::Stream& stream); - -template void computeTrackExtensionResultsHandler<11>(const TrackITSExt* tracks, - const TrackExtensionCandidate<11>* candidates, - const int* candidateOffsets, - TrackExtensionResult<11>* results, - const TrackingFrameInfo** trackingFrameInfo, - const std::array layerxX0, - const int nTracks, - const int nLayers, - const float bz, - const float maxChi2ClusterAttachment, - const float maxChi2NDF, - const o2::base::Propagator* propagator, - const o2::base::PropagatorF::MatCorrType matCorrType, - const bool shiftRefToCluster, - gpu::Stream& stream); - template void countTrackletsInROFsHandler<11>(const IndexTableUtils<11>* utils, const ROFMaskTable<11>::View& rofMask, const int transitionId, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h index 4e8d5bcfea42a..a7b44e91d9093 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h @@ -113,14 +113,17 @@ GPUhdi() void IndexTableUtils::print() const } template -GPUhdi() int4 getBinsRect(const Cluster& currentCluster, const int layerIndex, - const float z1, const float z2, const float maxdeltaz, const float maxdeltaphi, +GPUhdi() int4 getBinsRect(const int layerIndex, + const float phi, + const float z, + const float maxdeltaz, + const float maxdeltaphi, const IndexTableUtils& utils) { - const float zRangeMin = o2::gpu::GPUCommonMath::Min(z1, z2) - maxdeltaz; - const float phiRangeMin = (maxdeltaphi > o2::constants::math::PI) ? 0.f : currentCluster.phi - maxdeltaphi; - const float zRangeMax = o2::gpu::GPUCommonMath::Max(z1, z2) + maxdeltaz; - const float phiRangeMax = (maxdeltaphi > o2::constants::math::PI) ? o2::constants::math::TwoPI : currentCluster.phi + maxdeltaphi; + const float zRangeMin = z - maxdeltaz; + const float phiRangeMin = (maxdeltaphi > o2::constants::math::PI) ? 0.f : phi - maxdeltaphi; + const float zRangeMax = z + maxdeltaz; + const float phiRangeMax = (maxdeltaphi > o2::constants::math::PI) ? o2::constants::math::TwoPI : phi + maxdeltaphi; if (zRangeMax < -utils.getLayerZ(layerIndex) || zRangeMin > utils.getLayerZ(layerIndex) || zRangeMin > zRangeMax) { @@ -133,5 +136,15 @@ GPUhdi() int4 getBinsRect(const Cluster& currentCluster, const int layerIndex, utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))}; } +template +GPUhdi() int4 getBinsRect(const Cluster& currentCluster, const int layerIndex, + const float z1, const float z2, const float maxdeltaz, const float maxdeltaphi, + const IndexTableUtils& utils) +{ + const float zMean = 0.5f * (z1 + z2); + const float zDelta = 0.5f * o2::gpu::GPUCommonMath::Abs(z1 - z2) + maxdeltaz; + return getBinsRect(layerIndex, currentCluster.phi, zMean, zDelta, maxdeltaphi, utils); +} + } // namespace o2::its #endif /* TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h index a8e2c37e261fb..1584b5dcd79aa 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/ROFLookupTables.h @@ -37,6 +37,7 @@ namespace o2::its // Layer timing definition struct LayerTiming { using BCType = TimeStampType; + using BCRange = dataformats::RangeReference; BCType mNROFsTF{0}; // number of ROFs per timeframe BCType mROFLength{0}; // ROF length in BC BCType mROFDelay{0}; // delay of ROFs wrt start of first orbit in TF in BC @@ -110,26 +111,31 @@ struct LayerTiming { } // return clamped ROF range with strictly positive overlap with timestamp interval - GPUhdi() int2 getROFRange(TimeStamp ts) const noexcept + GPUhdi() BCRange getROFRange(TimeStamp ts) const noexcept { - if (mNROFsTF == 0) { - return {1, 0}; - } - const float lower = ts.getTimeStamp() - ts.getTimeStampError(); const float upper = ts.getTimeStamp() + ts.getTimeStampError(); - const int maxROF = static_cast(mNROFsTF) - 1; - int2 range{ - o2::gpu::CAMath::Clamp(static_cast(getROF(lower - mROFAddTimeErr)), 0, maxROF), - o2::gpu::CAMath::Clamp(static_cast(getROF(upper + mROFAddTimeErr)), 0, maxROF)}; + return getROFRange(lower, upper); + } + + GPUhdi() BCRange getROFRange(TimeEstBC ts) const noexcept + { + return getROFRange(static_cast(ts.lower()), static_cast(ts.upper())); + } + + GPUhdi() BCRange getROFRange(float lower, float upper) const noexcept + { + const BCType maxROF = mNROFsTF - 1; + BCType first = o2::gpu::CAMath::Clamp(getROF(lower - mROFAddTimeErr), BCType{0}, maxROF); + BCType last = o2::gpu::CAMath::Clamp(getROF(upper + mROFAddTimeErr), BCType{0}, maxROF); - if (range.x <= range.y && !intersectROF(static_cast(range.x), lower, upper)) { - ++range.x; + if (first <= last && !intersectROF(first, lower, upper)) { + ++first; } - if (range.y >= range.x && !intersectROF(static_cast(range.y), lower, upper)) { - --range.y; + if (last >= first && !intersectROF(last, lower, upper)) { + --last; } - return range; + return {first, first <= last ? static_cast(last - first + 1) : BCType{0}}; } #ifndef GPUCA_GPUCODE diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index 3fef2dc640cbc..9a31e4014bff5 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -212,6 +212,10 @@ struct TimeFrame { virtual size_t getNumberOfNeighbours() const; size_t getNumberOfTracks() const; size_t getNumberOfUsedClusters() const; + void resetTrackExtensionCounters(); + void addTrackExtensionCounters(size_t nTracks, size_t nClusters); + size_t getNExtendedTracks() const { return mNExtendedTracks; } + size_t getNExtendedClusters() const { return mNExtendedClusters; } /// memory management void setMemoryPool(std::shared_ptr pool); @@ -280,6 +284,8 @@ struct TimeFrame { std::vector> mCells; bounded_vector mTracks; bounded_vector mTracksLabel; + size_t mNExtendedTracks = 0; + size_t mNExtendedClusters = 0; std::vector> mCellsNeighbours; std::vector> mCellsNeighboursTopology; std::vector> mCellsLookupTable; @@ -604,6 +610,20 @@ inline size_t TimeFrame::getNumberOfUsedClusters() const return nClusters; } +template +inline void TimeFrame::resetTrackExtensionCounters() +{ + mNExtendedTracks = 0; + mNExtendedClusters = 0; +} + +template +inline void TimeFrame::addTrackExtensionCounters(size_t nTracks, size_t nClusters) +{ + mNExtendedTracks += nTracks; + mNExtendedClusters += nClusters; +} + } // namespace its } // namespace o2 diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h deleted file mode 100644 index 5ff5bc4c0828b..0000000000000 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h +++ /dev/null @@ -1,134 +0,0 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. -// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. -// All rights not expressly granted are reserved. -// -// This software is distributed under the terms of the GNU General Public -// License v3 (GPL Version 3), copied verbatim in the file "COPYING". -// -// In applying this license CERN does not waive the privileges and immunities -// granted to it by virtue of its status as an Intergovernmental Organization -// or submit itself to any jurisdiction. - -#ifndef TRACKINGITSU_INCLUDE_TRACKEXTENSIONCANDIDATE_H_ -#define TRACKINGITSU_INCLUDE_TRACKEXTENSIONCANDIDATE_H_ - -#include -#include - -#include "GPUCommonDef.h" -#include "DataFormatsITS/TrackITS.h" -#include "DataFormatsITS/TimeEstBC.h" -#include "ITStracking/Constants.h" -#include "ReconstructionDataFormats/Track.h" - -namespace o2::its -{ - -inline constexpr unsigned int kExtendedPatternShift = 24; -inline constexpr int kMaxLayersInTrackPattern = 8; - -template -GPUhdi() constexpr uint32_t makeAddedClustersPatternMask() -{ - return (NLayers >= 32) ? 0xffffffffu : ((1u << NLayers) - 1u); -} - -template -GPUhdi() void applyExtendedClustersPattern(TrackITSExt& track, uint32_t diff) -{ - diff &= makeAddedClustersPatternMask(); - track.setUserField(static_cast(diff)); - if constexpr (NLayers <= kMaxLayersInTrackPattern) { - track.setPattern(track.getPattern() | (diff << kExtendedPatternShift)); - } else { - (void)track; - } -} - -template -GPUhdi() uint32_t getAddedClustersPattern(const TrackITSExt& track) -{ - const auto mask = makeAddedClustersPatternMask(); - if constexpr (NLayers <= kMaxLayersInTrackPattern) { - const auto diff = (track.getPattern() >> kExtendedPatternShift) & mask; - if (diff) { - return diff; - } - } - return track.getUserField() & mask; -} - -GPUhdi() void clearAddedClustersPattern(TrackITSExt& track) -{ - track.setUserField(0); - track.getParamOut().setUserField(0); -} - -template -struct TrackExtensionHypothesis { - o2::track::TrackParCov param; - std::array clusters{}; - TimeStamp time; - float chi2{0.f}; - int nClusters{0}; - int edgeLayer{constants::UnusedIndex}; -}; - -template -struct TrackExtensionCandidate { - static constexpr float InvalidChi2 = 1.e20f; - - GPUhdi() TrackExtensionCandidate() { reset(); } - - GPUhdi() void reset() - { - trackIndex = -1; - nAddedClusters = 0; - resultIndex = -1; - chi2 = InvalidChi2; - for (int iLayer{0}; iLayer < NLayers; ++iLayer) { - addedClusters[iLayer] = constants::UnusedIndex; - } - } - - GPUhdi() bool isValidForTrack(int index) const - { - return trackIndex == index && nAddedClusters > 0; - } - - int trackIndex{-1}; - std::array addedClusters; - int nAddedClusters{0}; - int resultIndex{-1}; - float chi2{InvalidChi2}; -}; - -template -GPUhdi() bool isBetterTrackExtensionCandidate(const TrackExtensionCandidate& a, const TrackExtensionCandidate& b) -{ - return (a.nAddedClusters > b.nAddedClusters) || (a.nAddedClusters == b.nAddedClusters && a.chi2 < b.chi2); -} - -template -struct TrackExtensionResult { - GPUhdi() void reset() - { - candidate.reset(); - } - - GPUhdi() bool isValid() const { return candidate.trackIndex >= 0 && candidate.nAddedClusters > 0; } - - TrackExtensionCandidate candidate; - TrackITSExt track; -}; - -inline constexpr int MaxTrackExtensionCandidatesPerTrack = 4; - -inline constexpr size_t getFlatTrackExtensionCandidateIndex(size_t trackIndex, size_t candidateIndex) -{ - return trackIndex * MaxTrackExtensionCandidatesPerTrack + candidateIndex; -} - -} // namespace o2::its - -#endif /* TRACKINGITSU_INCLUDE_TRACKEXTENSIONCANDIDATE_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionHypothesis.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionHypothesis.h new file mode 100644 index 0000000000000..ddfc88c239f36 --- /dev/null +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionHypothesis.h @@ -0,0 +1,56 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#ifndef TRACKINGITSU_INCLUDE_TRACKEXTENSIONHYPOTHESIS_H_ +#define TRACKINGITSU_INCLUDE_TRACKEXTENSIONHYPOTHESIS_H_ + +#include + +#include "GPUCommonDef.h" +#include "DataFormatsITS/TimeEstBC.h" +#include "ITStracking/Constants.h" +#include "ITStracking/TrackITSInternal.h" +#include "ReconstructionDataFormats/Track.h" + +namespace o2::its +{ + +template +struct TrackExtensionHypothesis { + TrackExtensionHypothesis() = default; + GPUhdi() TrackExtensionHypothesis(const TrackITSInternal& track, bool outward) + { + initialiseFromTrack(track, outward); + } + + GPUhdi() void initialiseFromTrack(const TrackITSInternal& track, bool outward) + { + param = outward ? track.paramOut : track.paramIn; + time = track.time; + chi2 = track.getChi2(); + nClusters = track.getNClusters(); + edgeLayer = outward ? track.getLastClusterLayer() : track.getFirstClusterLayer(); + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + clusters[iLayer] = track.getClusterIndex(iLayer); + } + } + + o2::track::TrackParCov param; + std::array clusters{}; + TimeEstBC time; + float chi2{0.f}; + int nClusters{0}; + int edgeLayer{constants::UnusedIndex}; +}; + +} // namespace o2::its + +#endif /* TRACKINGITSU_INCLUDE_TRACKEXTENSIONHYPOTHESIS_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h index 8cd20262edf14..76365b6ebf466 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h @@ -17,7 +17,6 @@ #include "GPUCommonDef.h" #include "GPUCommonMath.h" -#include "CommonConstants/MathConstants.h" #include "DetectorsBase/Propagator.h" #include "ITStracking/Cluster.h" @@ -25,17 +24,12 @@ #include "ITStracking/IndexTableUtils.h" #include "ITStracking/MathUtils.h" #include "ITStracking/ROFLookupTables.h" -#include "ITStracking/TrackExtensionCandidate.h" +#include "ITStracking/TrackExtensionHypothesis.h" +#include "ITStracking/TrackHelpers.h" namespace o2::its { -template -GPUhdi() bool isBetterTrackExtensionHypothesis(const TrackExtensionHypothesis& a, const TrackExtensionHypothesis& b) -{ - return (a.nClusters > b.nClusters) || (a.nClusters == b.nClusters && a.chi2 < b.chi2); -} - template GPUhdi() void addTrackExtensionHypothesisToBeam(const TrackExtensionHypothesis& hypo, TrackExtensionHypothesis* beam, @@ -49,77 +43,37 @@ GPUhdi() void addTrackExtensionHypothesisToBeam(const TrackExtensionHypothesis -GPUhdi() int4 getTrackExtensionBinsAt(const IndexTableUtils& utils, - const int layer, - const float phi, - const float deltaPhi, - const float z, - const float deltaZ) -{ - const float zRangeMin = z - deltaZ; - const float zRangeMax = z + deltaZ; - if (zRangeMax < -utils.getLayerZ(layer) || zRangeMin > utils.getLayerZ(layer) || zRangeMin > zRangeMax) { - return {-1, -1, -1, -1}; - } - const float phiRangeMin = (deltaPhi > o2::constants::math::PI) ? 0.f : phi - deltaPhi; - const float phiRangeMax = (deltaPhi > o2::constants::math::PI) ? o2::constants::math::TwoPI : phi + deltaPhi; - return {o2::gpu::CAMath::Max(0, utils.getZBinIndex(layer, zRangeMin)), - utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMin)), - o2::gpu::CAMath::Min(utils.getNzBins() - 1, utils.getZBinIndex(layer, zRangeMax)), - utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))}; -} - -template -GPUhdi() int getTrackExtensionFirstClusterLayer(const TrackITSExt& track) +GPUhdi() void updateTrackFromExtensionHypothesis(const TrackExtensionHypothesis& hypo, + const bool outward, + const int nLayers, + TrackITSInternal& track) { - const uint32_t pattern = track.getPattern(); - for (int iLayer{0}; iLayer < NLayers; ++iLayer) { - if (pattern & (0x1u << iLayer)) { - return iLayer; - } + if (outward) { + track.paramOut = hypo.param; + } else { + track.paramIn = hypo.param; } - return constants::UnusedIndex; -} - -template -GPUhdi() int getTrackExtensionLastClusterLayer(const TrackITSExt& track) -{ - const uint32_t pattern = track.getPattern(); - for (int iLayer{NLayers}; iLayer-- > 0;) { - if (pattern & (0x1u << iLayer)) { - return iLayer; + track.time = hypo.time; + track.setChi2(hypo.chi2); + for (int iLayer{0}; iLayer < nLayers; ++iLayer) { + if (track.getClusterIndex(iLayer) == constants::UnusedIndex && hypo.clusters[iLayer] != constants::UnusedIndex) { + track.setClusterIndex(iLayer, hypo.clusters[iLayer]); } } - return constants::UnusedIndex; } template -GPUhdi() void initialiseTrackExtensionHypothesis(const TrackITSExt& track, - const bool outward, - TrackExtensionHypothesis& hypo) -{ - hypo.param = outward ? track.getParamOut() : track.getParamIn(); - hypo.time = track.getTimeStamp(); - hypo.chi2 = track.getChi2(); - hypo.nClusters = track.getNClusters(); - hypo.edgeLayer = outward ? getTrackExtensionLastClusterLayer(track) : getTrackExtensionFirstClusterLayer(track); - for (int iLayer{0}; iLayer < NLayers; ++iLayer) { - hypo.clusters[iLayer] = track.getClusterIndex(iLayer); - } -} - -template -GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, +GPUhdi() bool followTrackExtensionDirection(const TrackExtensionHypothesis& startHypothesis, const IndexTableUtils& utils, const typename ROFMaskTable::View& rofMask, const typename ROFOverlapTable::View& rofOverlaps, @@ -143,14 +97,14 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, const o2::base::PropagatorF::MatCorrType matCorrType, TrackExtensionHypothesis* activeHypotheses, TrackExtensionHypothesis* nextHypotheses, - TrackITSExt& updatedTrack) + TrackExtensionHypothesis& bestHypothesis) { const int step = outward ? 1 : -1; const int end = outward ? nLayers - 1 : 0; const int beamWidth = o2::gpu::CAMath::Max(beamWidthConfig, 1); int nActive{1}; int nNext{0}; - initialiseTrackExtensionHypothesis(track, outward, activeHypotheses[0]); + activeHypotheses[0] = startHypothesis; const int tableSize = utils.getNphiBins() * utils.getNzBins() + 1; for (int iLayer = activeHypotheses[0].edgeLayer + step; nActive > 0; iLayer += step) { @@ -177,12 +131,7 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, const float ePhi{o2::gpu::CAMath::Sqrt(hypo.param.getSigmaSnp2() / hypo.param.getCsp2())}; const float eZ{o2::gpu::CAMath::Sqrt(hypo.param.getSigmaZ2())}; - const int4 selectedBins = getTrackExtensionBinsAt(utils, - iLayer, - hypo.param.getPhi(), - nSigmaCutPhi * ePhi, - hypo.param.getZ(), - nSigmaCutZ * eZ); + const int4 selectedBins = getBinsRect(iLayer, hypo.param.getPhi(), hypo.param.getZ(), nSigmaCutZ * eZ, nSigmaCutPhi * ePhi, utils); if (selectedBins.x < 0) { continue; } @@ -193,7 +142,7 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, } const auto rofRange = rofOverlaps.getLayer(iLayer).getROFRange(hypo.time); - for (int rof = rofRange.x; rof <= rofRange.y; ++rof) { + for (int rof = rofRange.getFirstEntry(); rof < rofRange.getEntriesBound(); ++rof) { if (!rofMask.isROFEnabled(iLayer, rof)) { continue; } @@ -241,12 +190,7 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, updated.clusters[iLayer] = nextCluster.clusterId; ++updated.nClusters; updated.edgeLayer = iLayer; - const auto rofTS = rofOverlaps.getLayer(iLayer).getROFTimeBounds(rof, true); - const auto& ts = updated.time; - const float lower = o2::gpu::CAMath::Max(ts.getTimeStamp() - ts.getTimeStampError(), static_cast(rofTS.lower())); - const float upper = o2::gpu::CAMath::Min(ts.getTimeStamp() + ts.getTimeStampError(), static_cast(rofTS.upper())); - updated.time.setTimeStamp(0.5f * (lower + upper)); - updated.time.setTimeStampError(0.5f * (upper - lower)); + updated.time += rofOverlaps.getLayer(iLayer).getROFTimeBounds(rof, true); addTrackExtensionHypothesisToBeam(updated, nextHypotheses, nNext, beamWidth); } } @@ -265,14 +209,14 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, const TrackExtensionHypothesis* bestHypo{nullptr}; for (int iHypo{0}; iHypo < nActive; ++iHypo) { const auto& hypo = activeHypotheses[iHypo]; - if (hypo.nClusters == track.getNClusters()) { + if (hypo.nClusters == startHypothesis.nClusters) { continue; } const float maxChi2 = maxChi2NDF * static_cast(hypo.nClusters * 2 - 5); if (hypo.chi2 >= maxChi2) { continue; } - if (!bestHypo || isBetterTrackExtensionHypothesis(hypo, *bestHypo)) { + if (!bestHypo || track::isBetter(hypo.nClusters, hypo.chi2, bestHypo->nClusters, bestHypo->chi2)) { bestHypo = &hypo; } } @@ -280,19 +224,7 @@ GPUhdi() bool followTrackExtensionDirection(const TrackITSExt& track, return false; } - updatedTrack = track; - if (outward) { - updatedTrack.getParamOut() = bestHypo->param; - } else { - updatedTrack.getParamIn() = bestHypo->param; - } - updatedTrack.getTimeStamp() = bestHypo->time; - updatedTrack.setChi2(bestHypo->chi2); - for (int iLayer{0}; iLayer < nLayers; ++iLayer) { - if (updatedTrack.getClusterIndex(iLayer) == constants::UnusedIndex && bestHypo->clusters[iLayer] != constants::UnusedIndex) { - updatedTrack.setExternalClusterIndex(iLayer, bestHypo->clusters[iLayer], true); - } - } + bestHypothesis = *bestHypo; return true; } diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h index d244b39ff9d11..4077e57f7d571 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h @@ -16,25 +16,33 @@ #ifndef O2_ITS_TRACKING_TRACKHELPERS_H_ #define O2_ITS_TRACKING_TRACKHELPERS_H_ +#include "CommonConstants/MathConstants.h" #include "DataFormatsITS/TrackITS.h" #include "ITStracking/Cell.h" #include "ITStracking/Cluster.h" #include "ITStracking/Constants.h" #include "ITStracking/MathUtils.h" +#include "ITStracking/TrackITSInternal.h" #include "DetectorsBase/Propagator.h" #include "ReconstructionDataFormats/Track.h" namespace o2::its::track { -// Prefer 1) longer track 2) sorted in chi2 +GPUhdi() bool isBetter(const int nClustersA, const float chi2A, const int nClustersB, const float chi2B) +{ + return (nClustersA > nClustersB) || (nClustersA == nClustersB && chi2A < chi2B); +} + GPUhdi() bool isBetter(const o2::its::TrackITS& a, const o2::its::TrackITS& b) { - const auto ncla = a.getNumberOfClusters(); - const auto nclb = b.getNumberOfClusters(); - // is a as long as b ? then decide on chi2 - // otherwise prefer longer - return (ncla == nclb) ? (a.getChi2() < b.getChi2()) : ncla > nclb; + 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()); } // Find the populated interior layer closest to the radial midpoint. @@ -58,7 +66,7 @@ GPUdi() int selectReseedMidLayer(int minLayer, int maxLayer, const float* layerR return midLayer; } -GPUdi() void resetTrackCovariance(TrackITSExt& track) +GPUdi() void resetTrackCovariance(o2::track::TrackParCov& track) { track.resetCovariance(); track.setCov(track.getQ2Pt() * track.getQ2Pt() * track.getCov()[o2::track::CovLabels::kSigQ2Pt2], o2::track::CovLabels::kSigQ2Pt2); @@ -97,19 +105,20 @@ GPUdi() o2::track::TrackParCov buildTrackSeed(const Cluster& cluster1, } template -GPUdi() TrackITSExt seedTrackForRefit(const TrackSeed& seed, - const TrackingFrameInfo* const* foundTrackingFrameInfo, - const Cluster* const* unsortedClusters, - const float* layerRadii, - const float bz, - const int reseedIfShorter) +GPUdi() TrackITSInternal seedTrackForRefit(const TrackSeed& seed, + const TrackingFrameInfo* const* foundTrackingFrameInfo, + const Cluster* const* unsortedClusters, + const float* layerRadii, + const float bz, + const int reseedIfShorter) { - TrackITSExt temporaryTrack(seed); + TrackITSInternal temporaryTrack; + temporaryTrack.paramIn = static_cast(seed); int lrMin = NLayers; int lrMax = 0; for (int iL{0}; iL < NLayers; ++iL) { const int idx = seed.getCluster(iL); - temporaryTrack.setExternalClusterIndex(iL, idx, idx != constants::UnusedIndex); + temporaryTrack.setClusterIndex(iL, idx); if (idx != constants::UnusedIndex) { lrMin = o2::gpu::CAMath::Min(lrMin, iL); lrMax = o2::gpu::CAMath::Max(lrMax, iL); @@ -123,15 +132,17 @@ GPUdi() TrackITSExt seedTrackForRefit(const TrackSeed& seed, const auto& cluster0TF = foundTrackingFrameInfo[lrMin][seed.getCluster(lrMin)]; const auto& cluster1GL = unsortedClusters[lrMid][seed.getCluster(lrMid)]; const auto& cluster2GL = unsortedClusters[lrMax][seed.getCluster(lrMax)]; - temporaryTrack.getParamIn() = buildTrackSeed(cluster2GL, cluster1GL, cluster0TF, bz, true); + temporaryTrack.paramIn = buildTrackSeed(cluster2GL, cluster1GL, cluster0TF, bz, true); } } - resetTrackCovariance(temporaryTrack); + resetTrackCovariance(temporaryTrack.paramIn); return temporaryTrack; } -GPUdi() bool fitTrack(TrackITSExt& trk, +template +GPUdi() bool fitTrack(TrackITSInternal& trk, + o2::track::TrackParCov& param, int start, int end, int step, @@ -154,43 +165,43 @@ GPUdi() bool fitTrack(TrackITSExt& trk, const TrackingFrameInfo& trackingHit = tfInfos[iLayer][trk.getClusterIndex(iLayer)]; if (linRef) { - if (!trk.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame, *linRef, bz)) { + if (!param.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame, *linRef, bz)) { return false; } - if (!propagator->propagateToX(trk, *linRef, trackingHit.xTrackingFrame, bz, + if (!propagator->propagateToX(param, *linRef, trackingHit.xTrackingFrame, bz, o2::base::PropagatorImpl::MAX_SIN_PHI, o2::base::PropagatorImpl::MAX_STEP, matCorrType)) { return false; } if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE) { - if (!trk.correctForMaterial(*linRef, layerxX0[iLayer], layerxX0[iLayer] * constants::Radl * constants::Rho, true)) { + if (!param.correctForMaterial(*linRef, layerxX0[iLayer], layerxX0[iLayer] * constants::Radl * constants::Rho, true)) { continue; } } } else { - if (!trk.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame)) { + if (!param.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame)) { return false; } - if (!propagator->propagateToX(trk, trackingHit.xTrackingFrame, bz, + if (!propagator->propagateToX(param, trackingHit.xTrackingFrame, bz, o2::base::PropagatorImpl::MAX_SIN_PHI, o2::base::PropagatorImpl::MAX_STEP, matCorrType)) { return false; } if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE) { - if (!trk.correctForMaterial(layerxX0[iLayer], layerxX0[iLayer] * constants::Radl * constants::Rho, true)) { + if (!param.correctForMaterial(layerxX0[iLayer], layerxX0[iLayer] * constants::Radl * constants::Rho, true)) { continue; } } } - const auto predChi2{trk.getPredictedChi2Quiet(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)}; + const auto predChi2{param.getPredictedChi2Quiet(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)}; if ((nCl >= 3 && predChi2 > chi2clcut) || predChi2 < 0.f) { return false; } trk.setChi2(trk.getChi2() + predChi2); - if (!trk.o2::track::TrackParCov::update(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)) { + if (!param.o2::track::TrackParCov::update(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)) { return false; } if (linRef && shiftRefToCluster) { @@ -200,36 +211,30 @@ GPUdi() bool fitTrack(TrackITSExt& trk, nCl++; } - return o2::gpu::CAMath::Abs(trk.getQ2Pt()) < maxQoverPt && trk.getChi2() < chi2ndfcut * (float)((nCl * 2) - 5); + return o2::gpu::CAMath::Abs(param.getQ2Pt()) < maxQoverPt && trk.getChi2() < chi2ndfcut * (float)((nCl * 2) - 5); } template -GPUdi() bool refitTrack(const TrackSeed& trackSeed, - TrackITSExt& temporaryTrack, - float chi2clcut, - float chi2ndfcut, - const float bz, +GPUdi() bool refitTrack(TrackITSInternal& track, const TrackingFrameInfo* const* tfInfos, - const Cluster* const* clusters, const float* layerxX0, - const float* layerRadii, - const float* minPt, + const int nLayers, + const float bz, + const float chi2clcut, + const float chi2ndfcut, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - const int reseedIfShorter, const bool shiftRefToCluster, - const bool repeatRefitOut) + const bool repeatRefitOut, + const float minPt = -1.f) { - temporaryTrack = seedTrackForRefit(trackSeed, - tfInfos, - clusters, - layerRadii, - bz, - reseedIfShorter); - o2::track::TrackPar linRef{temporaryTrack}; - bool fitSuccess = fitTrack(temporaryTrack, + o2::track::TrackPar linRef{track.paramIn}; + resetTrackCovariance(track.paramIn); + track.setChi2(0); + bool fitSuccess = fitTrack(track, + track.paramIn, 0, - NLayers, + nLayers, 1, chi2clcut, chi2ndfcut, @@ -245,12 +250,14 @@ GPUdi() bool refitTrack(const TrackSeed& trackSeed, if (!fitSuccess) { return false; } - temporaryTrack.getParamOut() = temporaryTrack.getParamIn(); - linRef = temporaryTrack.getParamOut(); // use refitted track as lin.reference - resetTrackCovariance(temporaryTrack); - temporaryTrack.setChi2(0); - fitSuccess = fitTrack(temporaryTrack, - NLayers - 1, + + track.paramOut = track.paramIn; + linRef = track.paramOut; + resetTrackCovariance(track.paramIn); + track.setChi2(0); + fitSuccess = fitTrack(track, + track.paramIn, + nLayers - 1, -1, -1, chi2clcut, @@ -264,36 +271,81 @@ GPUdi() bool refitTrack(const TrackSeed& trackSeed, matCorrType, &linRef, shiftRefToCluster); - if (!fitSuccess || temporaryTrack.getPt() < minPt[NLayers - temporaryTrack.getNClusters()]) { + if (!fitSuccess) { + return false; + } + if (minPt > 0.f && track.getPt() < minPt) { return false; } if (repeatRefitOut) { // repeat outward refit seeding and linearizing with the stable inward fit result - o2::track::TrackParCov saveInw{temporaryTrack}; + o2::track::TrackParCov saveInw{track.paramIn}; linRef = saveInw; // use refitted track as lin.reference - float saveChi2 = temporaryTrack.getChi2(); - track::resetTrackCovariance(temporaryTrack); - temporaryTrack.setChi2(0); - fitSuccess = o2::its::track::fitTrack(temporaryTrack, - 0, - NLayers, - 1, - chi2clcut, - chi2ndfcut, - o2::constants::math::VeryBig, - 0, - bz, - tfInfos, - layerxX0, - propagator, - matCorrType, - &linRef, - shiftRefToCluster); + float saveChi2 = track.getChi2(); + track.paramOut = saveInw; + track::resetTrackCovariance(track.paramOut); + track.setChi2(0); + fitSuccess = fitTrack(track, + track.paramOut, + 0, + nLayers, + 1, + chi2clcut, + chi2ndfcut, + o2::constants::math::VeryBig, + 0, + bz, + tfInfos, + layerxX0, + propagator, + matCorrType, + &linRef, + shiftRefToCluster); if (!fitSuccess) { return false; } - temporaryTrack.getParamOut() = temporaryTrack.getParamIn(); - temporaryTrack.getParamIn() = saveInw; - temporaryTrack.setChi2(saveChi2); + track.paramIn = saveInw; + track.setChi2(saveChi2); + } + return true; +} + +template +GPUdi() bool refitTrack(const TrackSeed& trackSeed, + TrackITSInternal& temporaryTrack, + float chi2clcut, + float chi2ndfcut, + const float bz, + const TrackingFrameInfo* const* tfInfos, + const Cluster* const* clusters, + const float* layerxX0, + const float* layerRadii, + const float* minPt, + const o2::base::Propagator* propagator, + const o2::base::PropagatorF::MatCorrType matCorrType, + const int reseedIfShorter, + const bool shiftRefToCluster, + const bool repeatRefitOut) +{ + temporaryTrack = seedTrackForRefit(trackSeed, + tfInfos, + clusters, + layerRadii, + bz, + reseedIfShorter); + bool fitSuccess = refitTrack(temporaryTrack, + tfInfos, + layerxX0, + NLayers, + bz, + chi2clcut, + chi2ndfcut, + propagator, + matCorrType, + shiftRefToCluster, + repeatRefitOut, + minPt[NLayers - temporaryTrack.getNClusters()]); + if (!fitSuccess) { + return false; } return true; } diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackITSInternal.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackITSInternal.h new file mode 100644 index 0000000000000..1f56afb1468c8 --- /dev/null +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackITSInternal.h @@ -0,0 +1,113 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +#ifndef TRACKINGITSU_INCLUDE_TRACKITSINTERNAL_H_ +#define TRACKINGITSU_INCLUDE_TRACKITSINTERNAL_H_ + +#include + +#include "GPUCommonDef.h" +#include "DataFormatsITS/TrackITS.h" +#include "DataFormatsITS/TimeEstBC.h" +#include "ITStracking/Constants.h" +#include "ReconstructionDataFormats/Track.h" + +namespace o2::its +{ + +template +struct TrackITSInternal { + GPUhdi() TrackITSInternal() { resetClusters(); } + + GPUhdi() void resetClusters() + { + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + clusters[iLayer] = constants::UnusedIndex; + } + nClusters = 0; + } + + GPUhdi() int getClusterIndex(int layer) const { return clusters[layer]; } + + GPUhdi() void setClusterIndex(int layer, int cluster) + { + if (clusters[layer] == constants::UnusedIndex && cluster != constants::UnusedIndex) { + ++nClusters; + } else if (clusters[layer] != constants::UnusedIndex && cluster == constants::UnusedIndex) { + --nClusters; + } + clusters[layer] = cluster; + } + + GPUhdi() int getNClusters() const { return nClusters; } + GPUhdi() int getNumberOfClusters() const { return nClusters; } + GPUhdi() float getChi2() const { return chi2; } + GPUhdi() void setChi2(float value) { chi2 = value; } + GPUdi() float getPt() const { return paramIn.getPt(); } + + GPUhdi() uint32_t getPattern() const + { + uint32_t pattern{0}; + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + if (clusters[iLayer] != constants::UnusedIndex) { + pattern |= (0x1u << iLayer); + } + } + return pattern; + } + + GPUhdi() int getFirstClusterLayer() const + { + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + if (clusters[iLayer] != constants::UnusedIndex) { + return iLayer; + } + } + return constants::UnusedIndex; + } + + GPUhdi() int getLastClusterLayer() const + { + for (int iLayer{NLayers - 1}; iLayer >= 0; --iLayer) { + if (clusters[iLayer] != constants::UnusedIndex) { + return iLayer; + } + } + return constants::UnusedIndex; + } + + o2::track::TrackParCov paramIn; + o2::track::TrackParCov paramOut; + std::array clusters{}; + TimeEstBC time; + float chi2{0.f}; + int nClusters{0}; +}; + +template +GPUhdi() TrackITSExt makeTrackITSExt(const TrackITSInternal& track) +{ + TrackITSExt out; + out.getParamIn() = track.paramIn; + out.getParamOut() = track.paramOut; + out.setChi2(track.chi2); + out.getTimeStamp() = track.time.makeSymmetrical(); + for (int iLayer{0}; iLayer < NLayers; ++iLayer) { + if (track.clusters[iLayer] != constants::UnusedIndex) { + out.setExternalClusterIndex(iLayer, track.clusters[iLayer], true); + } + } + return out; +} + +} // namespace o2::its + +#endif /* TRACKINGITSU_INCLUDE_TRACKITSINTERNAL_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index 201ee0470d20b..aa28355c429a6 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -16,20 +16,17 @@ #ifndef TRACKINGITSU_INCLUDE_TRACKERTRAITS_H_ #define TRACKINGITSU_INCLUDE_TRACKERTRAITS_H_ -#include #include #include #include "DetectorsBase/Propagator.h" #include "ITStracking/Configuration.h" -#include "ITStracking/Constants.h" #include "ITStracking/IndexTableUtils.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/Cell.h" #include "ITStracking/BoundedAllocator.h" -#include "DataFormatsITS/TimeEstBC.h" -#include "ReconstructionDataFormats/Track.h" -#include "ITStracking/TrackExtensionCandidate.h" +#include "ITStracking/TrackExtensionHypothesis.h" +#include "ITStracking/TrackITSInternal.h" // #define OPTIMISATION_OUTPUT @@ -62,7 +59,7 @@ class TrackerTraits template void processNeighbours(int iteration, int defaultCellTopologyId, int iLevel, const bounded_vector& currentCellSeed, const bounded_vector& currentCellId, const bounded_vector& currentCellTopologyId, bounded_vector& updatedCellSeed, bounded_vector& updatedCellId, bounded_vector& updatedCellTopologyId); - void acceptTracks(int iteration, bounded_vector& tracks, bounded_vector>& firstClusters, size_t& nExtendedTracks, size_t& nExtendedClusters); + void acceptTracks(int iteration, bounded_vector& tracks, bounded_vector>& firstClusters); void markTracks(int iteration); void updateTrackingParameters(const std::vector& trkPars) @@ -93,8 +90,13 @@ class TrackerTraits protected: struct TrackFollowerScratch { - std::vector> activeHypotheses; - std::vector> nextHypotheses; + explicit TrackFollowerScratch(std::pmr::memory_resource* memoryResource) + : activeHypotheses(memoryResource), nextHypotheses(memoryResource) + { + } + + bounded_vector> activeHypotheses; + bounded_vector> nextHypotheses; }; bool finaliseTrackSeed(const TrackSeedN& seed, @@ -103,8 +105,7 @@ class TrackerTraits const TrackingFrameInfo* const* tfInfos, const Cluster* const* unsortedClusters, const o2::base::Propagator* propagator); - bool trackFollowing(TrackITSExt* track, bool outward, const int iteration, TrackFollowerScratch& scratch); - bool refitExtendedTrack(TrackITSExt& track, const int iteration); + bool trackFollowing(TrackITSInternal* track, bool outward, const int iteration, TrackFollowerScratch& scratch); o2::gpu::GPUChainITS* mChain = nullptr; TimeFrame* mTimeFrame; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h index d80974e90a4ac..fecdacc817e37 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h @@ -96,7 +96,8 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper TrackingMode::getTrackingParameters(TrackingMode if (trackParams.size() > 3 && tc.doUPCIteration) { trackParams[3].PassFlags.set(IterationStep::UseUPCMask, IterationStep::RebuildClusterLUT, IterationStep::SelectUPCVertices); } - float bFactor = std::abs(o2::base::Propagator::Instance()->getNominalBz()) / 5.0066791f; float bFactorTracklets = bFactor < 0.01f ? 1.f : bFactor; // for tracklets only @@ -220,12 +219,6 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode p.TrackFollowerNSigmaCutZ = tc.trackFollowerNSigmaCutZ; p.TrackFollowerNSigmaCutPhi = tc.trackFollowerNSigmaCutPhi; p.TrackFollowerBeamWidth = std::max(1, tc.trackFollowerBeamWidth); - if (tc.trackFollower & 0x1) { - p.PassFlags.set(IterationStep::TrackFollowerTop); - } - if (tc.trackFollower & 0x2) { - p.PassFlags.set(IterationStep::TrackFollowerBot); - } p.PrintMemory = tc.printMemory; p.MaxMemory = tc.maxMemory; @@ -240,6 +233,12 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode if (iter < constants::MaxIter) { p.MaxHoles = tc.maxHolesIter[iter]; p.HoleLayerMask = tc.holeLayerMaskIter[iter]; + if (tc.trackFollowerTop[iter]) { + p.PassFlags.set(IterationStep::TrackFollowerTop); + } + if (tc.trackFollowerBot[iter]) { + p.PassFlags.set(IterationStep::TrackFollowerBot); + } } if (tc.useMatCorrTGeo) { diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index 8375004cbfbad..232fc0178a30d 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -474,6 +474,7 @@ void TimeFrame::setFrameworkAllocator(ExternalAllocator* ext) template void TimeFrame::wipe() { + resetTrackExtensionCounters(); deepVectorClear(mTracks); deepVectorClear(mTracklets); deepVectorClear(mCells); diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index 57c99f2557840..95c2de41bfa97 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -75,6 +75,7 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e float timeFrame{0.}, timeTracklets{0.}, timeCells{0.}, timeNeighbours{0.}, timeRoads{0.}; size_t nTracklets{0}, nCells{0}, nNeighbours{0}; int nTracks{-static_cast(mTimeFrame->getNumberOfTracks())}; + mTimeFrame->resetTrackExtensionCounters(); iVertex = std::min(maxNvertices, 0); logger(std::format("==== ITS {} Tracking iteration {} summary ====", mTraits->getName(), iteration)); total += timeFrame = evaluateTask(&Tracker::initialiseTimeFrame, StateNames[mCurStep = TFInit], iteration, evalLog, iteration); @@ -92,6 +93,9 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e logger(std::format(" - Cell finding: {} cells found in {:.2f} ms", nCells, timeCells)); logger(std::format(" - Neighbours finding: {} neighbours found in {:.2f} ms", nNeighbours, timeNeighbours)); logger(std::format(" - Track finding: {} tracks found in {:.2f} ms", nTracks + mTimeFrame->getNumberOfTracks(), timeRoads)); + if (mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]) { + logger(std::format(" - Integrated track extension: {} tracks accepted using {} clusters", mTimeFrame->getNExtendedTracks(), mTimeFrame->getNExtendedClusters())); + } total += timeTracklets + timeCells + timeNeighbours + timeRoads; } } catch (const BoundedMemoryResource::MemoryLimitExceeded& err) { diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 7451fb3bff0a5..bc7157f4c53bf 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -672,8 +672,9 @@ bool TrackerTraits::finaliseTrackSeed(const TrackSeedN& seed, const Cluster* const* unsortedClusters, const o2::base::Propagator* propagator) { + TrackITSInternal internalTrack; if (!track::refitTrack(seed, - track, + internalTrack, mTrkParams[iteration].MaxChi2ClusterAttachment, mTrkParams[iteration].MaxChi2NDF, mBz, @@ -693,26 +694,39 @@ bool TrackerTraits::finaliseTrackSeed(const TrackSeedN& seed, const bool extendTop = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop]; const bool extendBot = mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]; if (!extendTop && !extendBot) { + track = makeTrackITSExt(internalTrack); return true; } - const auto backup = track; - auto best = track; - TrackFollowerScratch scratch; + const auto backup = internalTrack; + auto best = internalTrack; + uint32_t bestDiff{0}; + TrackFollowerScratch scratch{mMemoryPool.get()}; const uint32_t lastLayer = static_cast(mTrkParams[iteration].NLayers - 1); - auto finaliseExtensionCandidate = [&](TrackITSExt& candidate) { - const auto diff = (candidate.getPattern() & ~backup.getPattern()) & makeAddedClustersPatternMask(); - if (!diff || !refitExtendedTrack(candidate, iteration)) { + auto finaliseExtensionCandidate = [&](TrackITSInternal& candidate) { + const auto diff = (candidate.getPattern() & ~backup.getPattern()) & TrackITS::getLayerPatternMask(); + if (!diff || + !track::refitTrack(candidate, + tfInfos, + mTrkParams[iteration].LayerxX0.data(), + mTrkParams[iteration].NLayers, + mBz, + mTrkParams[iteration].MaxChi2ClusterAttachment, + mTrkParams[iteration].MaxChi2NDF, + propagator, + mTrkParams[iteration].CorrType, + mTrkParams[iteration].ShiftRefToCluster, + mTrkParams[iteration].RepeatRefitOut)) { return; } - applyExtendedClustersPattern(candidate, diff); if (track::isBetter(candidate, best)) { best = candidate; + bestDiff = diff; } }; - std::optional topResult, botResult; + std::optional> topResult, botResult; if (extendTop && backup.getLastClusterLayer() != lastLayer) { auto candidate = backup; if (trackFollowing(&candidate, true, iteration, scratch)) { @@ -742,7 +756,10 @@ bool TrackerTraits::finaliseTrackSeed(const TrackSeedN& seed, } } - track = best; + track = makeTrackITSExt(best); + if (bestDiff) { + track.setExtendedLayerPattern(bestDiff); + } return true; } @@ -758,7 +775,6 @@ void TrackerTraits::findRoads(const int iteration) tfInfos[iLayer] = mTimeFrame->getTrackingFrameInfoOnLayer(iLayer).data(); unsortedClusters[iLayer] = mTimeFrame->getUnsortedClusters()[iLayer].data(); } - size_t nExtendedTracks{0}, nExtendedClusters{0}; const auto topology = mTimeFrame->getTrackingTopologyView(); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { @@ -843,10 +859,7 @@ void TrackerTraits::findRoads(const int iteration) return track::isBetter(a, b); }); - acceptTracks(iteration, tracks, firstClusters, nExtendedTracks, nExtendedClusters); - } - if (mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop] || mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot]) { - LOGP(info, "Integrated track extension accepted {} tracks using {} clusters in iteration {}", nExtendedTracks, nExtendedClusters, iteration); + acceptTracks(iteration, tracks, firstClusters); } markTracks(iteration); } @@ -854,9 +867,7 @@ void TrackerTraits::findRoads(const int iteration) template void TrackerTraits::acceptTracks(int iteration, bounded_vector& tracks, - bounded_vector>& firstClusters, - size_t& nExtendedTracks, - size_t& nExtendedClusters) + bounded_vector>& firstClusters) { auto& trks = mTimeFrame->getTracks(); trks.reserve(trks.size() + tracks.size()); @@ -917,14 +928,15 @@ void TrackerTraits::acceptTracks(int iteration, if (track.getTimeStamp().getTimeStampError() > smallestROFHalf) { track.getTimeStamp().setTimeStampError(smallestROFHalf); } - const auto diff = getAddedClustersPattern(track); + const auto diff = track.getExtendedLayerPattern(); if (diff) { - ++nExtendedTracks; + size_t nExtendedClusters = 0; for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { nExtendedClusters += static_cast(diff & (0x1u << iLayer)); } + mTimeFrame->addTrackExtensionCounters(1, nExtendedClusters); } - clearAddedClustersPattern(track); + track.clearExtendedLayerPattern(); trks.emplace_back(track); if (mTrkParams[iteration].AllowSharingFirstCluster) { @@ -980,60 +992,7 @@ void TrackerTraits::markTracks(int iteration) } template -bool TrackerTraits::refitExtendedTrack(TrackITSExt& track, const int iteration) -{ - const auto propagator = o2::base::Propagator::Instance(); - const TrackingFrameInfo* tfInfos[NLayers]{}; - for (int iLayer = 0; iLayer < NLayers; ++iLayer) { - tfInfos[iLayer] = mTimeFrame->getTrackingFrameInfoOnLayer(iLayer).data(); - } - - o2::track::TrackPar linRef{track}; - track::resetTrackCovariance(track); - track.setChi2(0); - bool fitSuccess = track::fitTrack(track, - 0, - mTrkParams[iteration].NLayers, - 1, - mTrkParams[iteration].MaxChi2ClusterAttachment, - mTrkParams[iteration].MaxChi2NDF, - o2::constants::math::VeryBig, - 0, - mBz, - tfInfos, - mTrkParams[iteration].LayerxX0.data(), - propagator, - mTrkParams[iteration].CorrType, - &linRef, - mTrkParams[iteration].ShiftRefToCluster); - if (!fitSuccess) { - return false; - } - - track.getParamOut() = track.getParamIn(); - linRef = track.getParamOut(); - track::resetTrackCovariance(track); - track.setChi2(0); - fitSuccess = track::fitTrack(track, - mTrkParams[iteration].NLayers - 1, - -1, - -1, - mTrkParams[iteration].MaxChi2ClusterAttachment, - mTrkParams[iteration].MaxChi2NDF, - 50.f, - 0, - mBz, - tfInfos, - mTrkParams[iteration].LayerxX0.data(), - propagator, - mTrkParams[iteration].CorrType, - &linRef, - mTrkParams[iteration].ShiftRefToCluster); - return fitSuccess; -} - -template -bool TrackerTraits::trackFollowing(TrackITSExt* track, bool outward, const int iteration, TrackFollowerScratch& scratch) +bool TrackerTraits::trackFollowing(TrackITSInternal* track, bool outward, const int iteration, TrackFollowerScratch& scratch) { const int beamWidth = std::max(1, mTrkParams[iteration].TrackFollowerBeamWidth); if (static_cast(scratch.activeHypotheses.size()) < beamWidth) { @@ -1056,9 +1015,10 @@ bool TrackerTraits::trackFollowing(TrackITSExt* track, bool outward, co tfInfoPtrs[iLayer] = mTimeFrame->getTrackingFrameInfoOnLayer(iLayer).data(); } - TrackITSExt updated; + auto startHypothesis = TrackExtensionHypothesis{*track, outward}; + TrackExtensionHypothesis bestHypothesis; const bool ok = followTrackExtensionDirection( - *track, + startHypothesis, mTimeFrame->getIndexTableUtils(), mTimeFrame->getROFMaskView(), mTimeFrame->getROFOverlapTableView(), @@ -1082,20 +1042,12 @@ bool TrackerTraits::trackFollowing(TrackITSExt* track, bool outward, co mTrkParams[iteration].CorrType, scratch.activeHypotheses.data(), scratch.nextHypotheses.data(), - updated); + bestHypothesis); if (!ok) { return false; } - auto& trackParam = outward ? track->getParamOut() : track->getParamIn(); - trackParam = outward ? updated.getParamOut() : updated.getParamIn(); - track->setChi2(updated.getChi2()); - track->getTimeStamp() = updated.getTimeStamp(); - for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) { - if (track->getClusterIndex(iLayer) == constants::UnusedIndex && updated.getClusterIndex(iLayer) != constants::UnusedIndex) { - track->setExternalClusterIndex(iLayer, updated.getClusterIndex(iLayer), true); - } - } + updateTrackFromExtensionHypothesis(bestHypothesis, outward, mTrkParams[iteration].NLayers, *track); return true; } From d5d30e7b73e4cdd8a1486c50e0dd8550af12183e Mon Sep 17 00:00:00 2001 From: Maximiliano Puccio Date: Tue, 2 Jun 2026 14:06:12 +0200 Subject: [PATCH 4/4] ITS: rename away from beam for the track follower --- .../GPU/ITStrackingGPU/TimeFrameGPU.h | 2 +- .../GPU/ITStrackingGPU/TrackingKernels.h | 2 +- .../ITS/tracking/GPU/cuda/TimeFrameGPU.cu | 4 +-- .../tracking/GPU/cuda/TrackerTraitsGPU.cxx | 4 +-- .../ITS/tracking/GPU/cuda/TrackingKernels.cu | 24 +++++++-------- .../include/ITStracking/Configuration.h | 2 +- .../include/ITStracking/TrackFollower.h | 30 +++++++++---------- .../include/ITStracking/TrackingConfigParam.h | 6 ++-- .../ITSMFT/ITS/tracking/src/Configuration.cxx | 6 ++-- .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 12 ++++---- 10 files changed, 46 insertions(+), 46 deletions(-) diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index 3f574ff51a8c3..e2542d841c8bf 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -91,7 +91,7 @@ class TimeFrameGPU : public TimeFrame void createNeighboursDevice(const unsigned int layer); void createNeighboursLUTDevice(const int, const unsigned int); void createTrackITSExtDevice(const size_t); - void createTrackExtensionScratchDevice(const int nThreads, const int beamWidth); + void createTrackExtensionScratchDevice(const int nThreads, const int maxHypotheses); void downloadTrackITSExtDevice(); void downloadCellsNeighboursDevice(std::vector>&, const int); void downloadNeighboursLUTDevice(bounded_vector&, const int); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 4b3b36f513574..634ac5217a089 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -247,7 +247,7 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, const bool shiftRefToCluster, const int nLayers, const int phiBins, - const int beamWidth, + const int maxHypotheses, const bool extendTop, const bool extendBot, const float nSigmaCutPhi, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index 432ee62ab7b6e..bc2d90165ee96 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -583,10 +583,10 @@ void TimeFrameGPU::createTrackITSExtDevice(const size_t nSeeds) } template -void TimeFrameGPU::createTrackExtensionScratchDevice(const int nThreads, const int beamWidth) +void TimeFrameGPU::createTrackExtensionScratchDevice(const int nThreads, const int maxHypotheses) { GPUTimer timer("reserving track extension scratch"); - const size_t nHypotheses = static_cast(std::max(1, nThreads)) * std::max(1, beamWidth); + const size_t nHypotheses = static_cast(std::max(1, nThreads)) * std::max(1, maxHypotheses); GPULog("gpu-allocation: reserving {} track extension hypotheses per scratch buffer, for {:.2f} MB each.", nHypotheses, nHypotheses * sizeof(o2::its::TrackExtensionHypothesis) / constants::MB); mActiveTrackExtensionHypothesesDevice = nullptr; mNextTrackExtensionHypothesesDevice = nullptr; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 4cacf94684104..7ee18d69ceee0 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -371,7 +371,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) mTimeFrameGPU->getFrameworkAllocator()); mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size()); if (extendTracks) { - mTimeFrameGPU->createTrackExtensionScratchDevice(trackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerBeamWidth); + mTimeFrameGPU->createTrackExtensionScratchDevice(trackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerMaxHypotheses); } computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), @@ -400,7 +400,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[iteration].ShiftRefToCluster, this->mTrkParams[iteration].NLayers, this->mTrkParams[iteration].PhiBins, - this->mTrkParams[iteration].TrackFollowerBeamWidth, + this->mTrkParams[iteration].TrackFollowerMaxHypotheses, extendTop, extendBot, this->mTrkParams[iteration].TrackFollowerNSigmaCutPhi, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 69a2ff5be56d6..6d171da82f274 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -206,7 +206,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( const bool shiftRefToCluster, const int nLayers, const int phiBins, - const int beamWidthConfig, + const int maxHypothesesConfig, const bool extendTop, const bool extendBot, const float nSigmaCutPhi, @@ -236,10 +236,10 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( repeatRefitOut); if (refitSuccess) { if ((extendTop || extendBot) && activeHypothesesScratch && nextHypothesesScratch) { - const int beamWidth = o2::gpu::CAMath::Max(beamWidthConfig, 1); + const int maxHypotheses = o2::gpu::CAMath::Max(maxHypothesesConfig, 1); const int threadIndex = blockIdx.x * blockDim.x + threadIdx.x; - auto* activeHypotheses = activeHypothesesScratch + threadIndex * beamWidth; - auto* nextHypotheses = nextHypothesesScratch + threadIndex * beamWidth; + auto* activeHypotheses = activeHypothesesScratch + threadIndex * maxHypotheses; + auto* nextHypotheses = nextHypothesesScratch + threadIndex * maxHypotheses; const auto backupPattern = temporaryTrack.getPattern(); auto best = temporaryTrack; uint32_t bestDiff{0}; @@ -266,7 +266,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( layerxX0, nLayers, phiBins, - beamWidth, + maxHypotheses, bz, maxChi2ClusterAttachment, maxChi2NDF, @@ -301,7 +301,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( layerxX0, nLayers, phiBins, - beamWidth, + maxHypotheses, bz, maxChi2ClusterAttachment, maxChi2NDF, @@ -337,7 +337,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( layerxX0, nLayers, phiBins, - beamWidth, + maxHypotheses, bz, maxChi2ClusterAttachment, maxChi2NDF, @@ -370,7 +370,7 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( layerxX0, nLayers, phiBins, - beamWidth, + maxHypotheses, bz, maxChi2ClusterAttachment, maxChi2NDF, @@ -1343,7 +1343,7 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, const bool shiftRefToCluster, const int nLayers, const int phiBins, - const int beamWidth, + const int maxHypotheses, const bool extendTop, const bool extendBot, const float nSigmaCutPhi, @@ -1382,7 +1382,7 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, shiftRefToCluster, // bool nLayers, // int phiBins, // int - beamWidth, // int + maxHypotheses, // int extendTop, // bool extendBot, // bool nSigmaCutPhi, // float @@ -1584,7 +1584,7 @@ template void computeTrackSeedHandler(TrackSeed<7>* trackSeeds, const bool shiftRefToCluster, const int nLayers, const int phiBins, - const int beamWidth, + const int maxHypotheses, const bool extendTop, const bool extendBot, const float nSigmaCutPhi, @@ -1784,7 +1784,7 @@ template void computeTrackSeedHandler(TrackSeed<11>* trackSeeds, const bool shiftRefToCluster, const int nLayers, const int phiBins, - const int beamWidth, + const int maxHypotheses, const bool extendTop, const bool extendBot, const float nSigmaCutPhi, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 5c1dcf5216f51..ee8baf20e66c9 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -98,7 +98,7 @@ struct TrackingParameters { bool CreateArtefactLabels{false}; float TrackFollowerNSigmaCutZ = 1.f; float TrackFollowerNSigmaCutPhi = 1.f; - int TrackFollowerBeamWidth = 1; + int TrackFollowerMaxHypotheses = 1; bool PrintMemory = false; // print allocator usage in epilog report size_t MaxMemory = std::numeric_limits::max(); bool DropTFUponFailure = false; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h index 76365b6ebf466..4d3b197074909 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackFollower.h @@ -10,7 +10,7 @@ // or submit itself to any jurisdiction. /// \file TrackFollower.h -/// \brief Beam search used by CPU and GPU track extension. +/// \brief Hypothesis search used by CPU and GPU track extension. #ifndef TRACKINGITSU_INCLUDE_TRACKFOLLOWER_H_ #define TRACKINGITSU_INCLUDE_TRACKFOLLOWER_H_ @@ -31,24 +31,24 @@ namespace o2::its { template -GPUhdi() void addTrackExtensionHypothesisToBeam(const TrackExtensionHypothesis& hypo, - TrackExtensionHypothesis* beam, - int& nBeam, - const int beamWidth) +GPUhdi() void keepTrackExtensionHypothesis(const TrackExtensionHypothesis& hypo, + TrackExtensionHypothesis* keptHypotheses, + int& nKeptHypotheses, + const int maxHypotheses) { - if (nBeam < beamWidth) { - beam[nBeam++] = hypo; + if (nKeptHypotheses < maxHypotheses) { + keptHypotheses[nKeptHypotheses++] = hypo; return; } int worst{0}; - for (int i{1}; i < nBeam; ++i) { - if (track::isBetter(beam[worst].nClusters, beam[worst].chi2, beam[i].nClusters, beam[i].chi2)) { + for (int i{1}; i < nKeptHypotheses; ++i) { + if (track::isBetter(keptHypotheses[worst].nClusters, keptHypotheses[worst].chi2, keptHypotheses[i].nClusters, keptHypotheses[i].chi2)) { worst = i; } } - if (track::isBetter(hypo.nClusters, hypo.chi2, beam[worst].nClusters, beam[worst].chi2)) { - beam[worst] = hypo; + if (track::isBetter(hypo.nClusters, hypo.chi2, keptHypotheses[worst].nClusters, keptHypotheses[worst].chi2)) { + keptHypotheses[worst] = hypo; } } @@ -86,7 +86,7 @@ GPUhdi() bool followTrackExtensionDirection(const TrackExtensionHypothesis::max(); bool dropTFUponFailure = false; - bool fataliseUponFailure = true; // granular management of the fatalisation in async mode + bool fataliseUponFailure = true; // granular management of the fatalisation in async mode // Selections on tracks sharing clusters - bool allowSharingFirstCluster = false; // allow first cluster sharing among tracks + bool allowSharingFirstCluster = false; // allow first cluster sharing among tracks float sharedClusterMaxDeltaPhi = 0.05f; // Maximum allowed delta phi at the cluster position float sharedClusterMaxDeltaEta = 0.03f; // Maximum allowed delta eta at the cluster position bool sharedClusterOppositeSign = false; // Require opposite sign of the tracklets diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx index 096a9c7631f27..eb4fc399adc07 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx @@ -65,8 +65,8 @@ std::string TrackingParameters::asString() const top && bot ? "mix" : (top ? "top" : "bot"), TrackFollowerNSigmaCutZ, TrackFollowerNSigmaCutPhi); - if (TrackFollowerBeamWidth > 1) { - str += std::format(" Beam:{}", TrackFollowerBeamWidth); + if (TrackFollowerMaxHypotheses > 1) { + str += std::format(" MaxHypotheses:{}", TrackFollowerMaxHypotheses); } } if (std::numeric_limits::max() != MaxMemory) { @@ -218,7 +218,7 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode p.CreateArtefactLabels = tc.createArtefactLabels; p.TrackFollowerNSigmaCutZ = tc.trackFollowerNSigmaCutZ; p.TrackFollowerNSigmaCutPhi = tc.trackFollowerNSigmaCutPhi; - p.TrackFollowerBeamWidth = std::max(1, tc.trackFollowerBeamWidth); + p.TrackFollowerMaxHypotheses = std::max(1, tc.trackFollowerMaxHypotheses); p.PrintMemory = tc.printMemory; p.MaxMemory = tc.maxMemory; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index bc7157f4c53bf..93fcb78bb7379 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -994,12 +994,12 @@ void TrackerTraits::markTracks(int iteration) template bool TrackerTraits::trackFollowing(TrackITSInternal* track, bool outward, const int iteration, TrackFollowerScratch& scratch) { - const int beamWidth = std::max(1, mTrkParams[iteration].TrackFollowerBeamWidth); - if (static_cast(scratch.activeHypotheses.size()) < beamWidth) { - scratch.activeHypotheses.resize(beamWidth); + const int maxHypotheses = std::max(1, mTrkParams[iteration].TrackFollowerMaxHypotheses); + if (static_cast(scratch.activeHypotheses.size()) < maxHypotheses) { + scratch.activeHypotheses.resize(maxHypotheses); } - if (static_cast(scratch.nextHypotheses.size()) < beamWidth) { - scratch.nextHypotheses.resize(beamWidth); + if (static_cast(scratch.nextHypotheses.size()) < maxHypotheses) { + scratch.nextHypotheses.resize(maxHypotheses); } const Cluster* clustersPtrs[NLayers]{}; @@ -1031,7 +1031,7 @@ bool TrackerTraits::trackFollowing(TrackITSInternal* track, bo mTrkParams[iteration].LayerxX0.data(), mTrkParams[iteration].NLayers, mTrkParams[iteration].PhiBins, - beamWidth, + maxHypotheses, mBz, mTrkParams[iteration].MaxChi2ClusterAttachment, mTrkParams[iteration].MaxChi2NDF,