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/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 161283db2a2bc..ff541e0e5a839 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, @@ -208,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, @@ -222,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/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..43c45649b656a 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" @@ -301,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) { @@ -353,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, @@ -363,32 +368,53 @@ 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(); } + 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 571afe08fc209..55a0bc4d069e0 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,33 +110,388 @@ struct compare_track_chi2 { } }; -template +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) +{ + 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, + 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) 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], @@ -153,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; } } } @@ -584,6 +1078,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, @@ -1042,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, @@ -1058,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 @@ -1085,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) @@ -1106,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)); @@ -1131,6 +1760,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, @@ -1284,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, @@ -1297,26 +1971,87 @@ 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); /// 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, @@ -1470,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, @@ -1483,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/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/TrackExtensionCandidate.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h new file mode 100644 index 0000000000000..5ff5bc4c0828b --- /dev/null +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackExtensionCandidate.h @@ -0,0 +1,134 @@ +// 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/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..2362b6b2d9816 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h @@ -99,10 +99,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..201ee0470d20b 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -16,13 +16,20 @@ #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" // #define OPTIMISATION_OUTPUT @@ -55,7 +62,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); + 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) @@ -85,6 +92,20 @@ class TrackerTraits std::shared_ptr mTaskArena; protected: + 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); + 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..d80974e90a4ac 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/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index f17d961fc7bb7..57c99f2557840 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 diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index c4439dc74d29e..7451fb3bff0a5 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -14,9 +14,15 @@ /// #include +#include +#include #include +#include +#include +#include #include #include +#include #include #include @@ -30,6 +36,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" @@ -657,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) { @@ -669,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) { @@ -717,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); }); @@ -784,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()); @@ -851,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) { @@ -907,6 +979,126 @@ 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) +{ + 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) {