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) {