diff --git a/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h b/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h index b40c8feaffeba..f3eaedfa4cbea 100644 --- a/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h +++ b/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h @@ -92,8 +92,8 @@ class TrackITS : public o2::track::TrackParCov bool isBetter(const TrackITS& best, float maxChi2) const; - o2::track::TrackParCov& getParamIn() { return *this; } - const o2::track::TrackParCov& getParamIn() const { return *this; } + GPUhdi() o2::track::TrackParCov& getParamIn() { return *this; } + GPUhdi() const o2::track::TrackParCov& getParamIn() const { return *this; } GPUhdi() o2::track::TrackParCov& getParamOut() { return mParamOut; } GPUhdi() const o2::track::TrackParCov& getParamOut() const { return mParamOut; } diff --git a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrizationWithError.h b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrizationWithError.h index 2d11dee2147ca..42a58940897eb 100644 --- a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrizationWithError.h +++ b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrizationWithError.h @@ -46,7 +46,7 @@ class TrackParametrizationWithError : public TrackParametrization GPUd() TrackParametrizationWithError(const dim3_t& xyz, const dim3_t& pxpypz, const gpu::gpustd::array& cv, int sign, bool sectorAlpha = true, const PID pid = PID::Pion); - GPUdDefault() TrackParametrizationWithError(const TrackParametrizationWithError& src) = default; + GPUhdDefault() TrackParametrizationWithError(const TrackParametrizationWithError& src) = default; GPUdDefault() TrackParametrizationWithError(TrackParametrizationWithError&& src) = default; GPUhdDefault() TrackParametrizationWithError& operator=(const TrackParametrizationWithError& src) = default; GPUhdDefault() TrackParametrizationWithError& operator=(TrackParametrizationWithError&& src) = default; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index a30c129b4bccf..961be19287307 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -198,7 +198,9 @@ class TimeFrameGPU : public TimeFrame void loadTrackSeedsDevice(); void loadTrackSeedsChi2Device(); void loadRoadsDevice(); - void createTrackITSExtDevice(); + void loadTrackSeedsDevice(std::vector&); + void createTrackITSExtDevice(const std::vector& seeds); + void createTrackITSExtDevice(); // deprecated void downloadTrackITSExtDevice(); void initDeviceChunks(const int, const int); template @@ -232,6 +234,7 @@ class TimeFrameGPU : public TimeFrame Cluster** getDeviceArrayUnsortedClusters() const { return mUnsortedClustersDeviceArray; } Tracklet** getDeviceArrayTracklets() const { return mTrackletsDeviceArray; } CellSeed** getDeviceArrayCells() const { return mCellsDeviceArray; } + CellSeed* getDeviceTrackSeeds() { return mTrackSeedsDevice; } o2::track::TrackParCovF** getDeviceArrayTrackSeeds() { return mCellSeedsDeviceArray; } float** getDeviceArrayTrackSeedsChi2() { return mCellSeedsChi2DeviceArray; } void setDevicePropagator(const o2::base::PropagatorImpl*) override; @@ -263,6 +266,7 @@ class TimeFrameGPU : public TimeFrame std::array mTrackletsDevice; Tracklet** mTrackletsDeviceArray; std::array mCellsDevice; + CellSeed* mTrackSeedsDevice; CellSeed** mCellsDeviceArray; std::array mCellSeedsDevice; o2::track::TrackParCovF** mCellSeedsDeviceArray; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index 6e8f86d9d8d7a..b9ca1c19d60ff 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -575,6 +575,26 @@ void TimeFrameGPU::loadRoadsDevice() checkGPUError(cudaMemcpyAsync(mRoadsDevice, mRoads.data(), mRoads.size() * sizeof(Road), cudaMemcpyHostToDevice, mGpuStreams[0].get())); } +template +void TimeFrameGPU::loadTrackSeedsDevice(std::vector& seeds) +{ + LOGP(debug, "gpu-transfer: loading {} track seeds, for {} MB.", seeds.size(), seeds.size() * sizeof(CellSeed) / MB); + allocMemAsync(reinterpret_cast(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), &(mGpuStreams[0]), false); + checkGPUError(cudaHostRegister(seeds.data(), seeds.size() * sizeof(CellSeed), cudaHostRegisterPortable)); + checkGPUError(cudaMemcpyAsync(mTrackSeedsDevice, seeds.data(), seeds.size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0].get())); +} + +template +void TimeFrameGPU::createTrackITSExtDevice(const std::vector& seeds) +{ + mTrackITSExt.clear(); + mTrackITSExt.resize(seeds.size()); + LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB); + allocMemAsync(reinterpret_cast(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false); + checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get())); + checkGPUError(cudaHostRegister(mTrackITSExt.data(), seeds.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable)); +} + template void TimeFrameGPU::createTrackITSExtDevice() { @@ -589,8 +609,10 @@ void TimeFrameGPU::createTrackITSExtDevice() template void TimeFrameGPU::downloadTrackITSExtDevice() { - LOGP(debug, "gpu-transfer: downloading {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB); - checkGPUError(cudaMemcpyAsync(mTrackITSExt.data(), mTrackITSExtDevice, mRoads.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyDeviceToHost, mGpuStreams[0].get())); + LOGP(debug, "gpu-transfer: downloading {} tracks, for {} MB.", mTrackITSExt.size(), mTrackITSExt.size() * sizeof(o2::its::TrackITSExt) / MB); + checkGPUError(cudaMemcpyAsync(mTrackITSExt.data(), mTrackITSExtDevice, mTrackITSExt.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyDeviceToHost, mGpuStreams[0].get())); + checkGPUError(cudaHostUnregister(mTrackITSExt.data())); + discardResult(cudaDeviceSynchronize()); } template diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu index 694e84689305f..fa0593d9ce2dc 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu @@ -62,7 +62,7 @@ // #define GPUCA_CONSMEM const_cast(*gGPUConstantMemBuffer) // #endif // #define GPUCA_KRNL_BACKEND_CLASS GPUReconstructionHIPBackend -// // clang-format on +// clang-format on #endif // #include "GPUConstantMem.h" @@ -164,8 +164,7 @@ GPUd() bool fitTrack(TrackITSExt& track, float Bz, TrackingFrameInfo** tfInfos, const o2::base::Propagator* prop, - o2::base::PropagatorF::MatCorrType matCorrType = o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE, - bool debugPrint = false) + o2::base::PropagatorF::MatCorrType matCorrType = o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE) { for (int iLayer{start}; iLayer != end; iLayer += step) { if (track.getClusterIndex(iLayer) == constants::its::UnusedIndex) { @@ -672,6 +671,68 @@ GPUg() void computeLayerRoadsKernel( } } +template +GPUg() void fitTrackSeedsKernel( + CellSeed* trackSeeds, + TrackingFrameInfo** foundTrackingFrameInfo, + o2::its::TrackITSExt* tracks, + const size_t nSeeds, + const float Bz, + const int startLevel, + float maxChi2ClusterAttachment, + float maxChi2NDF, + const o2::base::Propagator* propagator) +{ + for (int iCurrentTrackSeedIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackSeedIndex < nSeeds; iCurrentTrackSeedIndex += blockDim.x * gridDim.x) { + auto& seed = trackSeeds[iCurrentTrackSeedIndex]; + if (seed.getQ2Pt() > 1.e3 || seed.getChi2() > maxChi2NDF * ((startLevel + 2) * 2 - (nLayers - 2))) { + continue; + } + TrackITSExt temporaryTrack{seed}; + temporaryTrack.resetCovariance(); + temporaryTrack.setChi2(0); + int* clusters = seed.getClusters(); + for (int iL{0}; iL < 7; ++iL) { + temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::its::UnusedIndex); + } + bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, + 0, // int lastLayer, + nLayers, // int firstLayer, + 1, // int firstCluster, + maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, + maxChi2NDF, // float maxChi2NDF, + o2::constants::math::VeryBig, // float maxQoverPt, + 0, // nCl, + Bz, // float Bz, + foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, + propagator, // const o2::base::Propagator* propagator, + o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE); // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT + if (!fitSuccess) { + continue; + } + temporaryTrack.getParamOut() = temporaryTrack.getParamIn(); + temporaryTrack.resetCovariance(); + temporaryTrack.setChi2(0); + + fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, + nLayers - 1, // int lastLayer, + -1, // int firstLayer, + -1, // int firstCluster, + maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, + maxChi2NDF, // float maxChi2NDF, + 50.f, // float maxQoverPt, + 0, // nCl, + Bz, // float Bz, + foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, + propagator, // const o2::base::Propagator* propagator, + o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE); // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT + if (!fitSuccess) { + continue; + } + tracks[iCurrentTrackSeedIndex] = temporaryTrack; + } +} + template GPUg() void fitTracksKernel( Cluster** foundClusters, @@ -689,7 +750,6 @@ GPUg() void fitTracksKernel( float maxChi2NDF, const o2::base::Propagator* propagator) { - o2::track::TrackParCovF track; for (int iCurrentRoadIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentRoadIndex < nRoads; iCurrentRoadIndex += blockDim.x * gridDim.x) { auto& currentRoad{roads[iCurrentRoadIndex]}; int clusters[nLayers]; @@ -747,65 +807,60 @@ GPUg() void fitTracksKernel( for (size_t iC = 0; iC < nLayers; ++iC) { temporaryTrack.setExternalClusterIndex(iC, clusters[iC], clusters[iC] != constants::its::UnusedIndex); } - bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, - lastCellLevel - 1, // int start, - -1, // int end, - -1, // int step, - maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, - maxChi2NDF, // float maxChi2NDF, - 1.e3, // float maxQoverPt, - 3, // int nCl, - Bz, // float Bz, - foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, - propagator, // const o2::base::Propagator* propagator, - o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE, // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT - iCurrentRoadIndex < 5); // Debug print + bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, + lastCellLevel - 1, // int start, + -1, // int end, + -1, // int step, + maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, + maxChi2NDF, // float maxChi2NDF, + 1.e3, // float maxQoverPt, + 3, // int nCl, + Bz, // float Bz, + foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, + propagator, // const o2::base::Propagator* propagator, + o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE); // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT + if (!fitSuccess) { - // printf("===> track %d died at %d iteration!\n", iCurrentRoadIndex, 1); continue; } temporaryTrack.resetCovariance(); temporaryTrack.setChi2(0); - fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, - 0, // int lastLayer, - 7, // int firstLayer, - 1, // int firstCluster, - maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, - maxChi2NDF, // float maxChi2NDF, - o2::constants::math::VeryBig, // float maxQoverPt, - 0, // nCl, - Bz, // float Bz, - foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, - propagator, // const o2::base::Propagator* propagator, - o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE, // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT - iCurrentRoadIndex < 5); // Debug print + fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, + 0, // int lastLayer, + 7, // int firstLayer, + 1, // int firstCluster, + maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, + maxChi2NDF, // float maxChi2NDF, + o2::constants::math::VeryBig, // float maxQoverPt, + 0, // nCl, + Bz, // float Bz, + foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, + propagator, // const o2::base::Propagator* propagator, + o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE); // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT if (!fitSuccess) { - // printf("===> track %d died at %d iteration!\n", iCurrentRoadIndex, 2); continue; } temporaryTrack.getParamOut() = temporaryTrack; temporaryTrack.resetCovariance(); temporaryTrack.setChi2(0); - fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, - 6 /* NL - 1 */, // int lastLayer, - -1, // int firstLayer, - -1, // int firstCluster, - maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, - maxChi2NDF, // float maxChi2NDF, - 50., // float maxQoverPt, - 0, // nCl, - Bz, // float Bz, - foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, - propagator, // const o2::base::Propagator* propagator, - o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE, // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT - iCurrentRoadIndex < 5); // Debug print + fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, + 6 /* NL - 1 */, // int lastLayer, + -1, // int firstLayer, + -1, // int firstCluster, + maxChi2ClusterAttachment, // float maxChi2ClusterAttachment, + maxChi2NDF, // float maxChi2NDF, + 50., // float maxQoverPt, + 0, // nCl, + Bz, // float Bz, + foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, + propagator, // const o2::base::Propagator* propagator, + o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE); // o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT + if (!fitSuccess) { - // printf("===> track %d died at %d iteration!\n", iCurrentRoadIndex, 3); continue; } tracks[iCurrentRoadIndex] = temporaryTrack; - // printf("===> track %d survived: ncl: %d -> ncl: %d\n", iCurrentRoadIndex, temporaryTrack.getNumberOfClusters(), tracks[iCurrentRoadIndex].getNumberOfClusters()); } } @@ -1060,16 +1115,24 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) } template -void TrackerTraitsGPU::findCellsNeighbours(const int iteration){}; +void TrackerTraitsGPU::findCellsNeighbours(const int iteration) +{ +} template -void TrackerTraitsGPU::findRoads(const int iteration){}; +void TrackerTraitsGPU::findRoads(const int iteration) +{ +} template -void TrackerTraitsGPU::findTracks(){}; +void TrackerTraitsGPU::findTracks() +{ +} template -void TrackerTraitsGPU::extendTracks(const int iteration){}; +void TrackerTraitsGPU::extendTracks(const int iteration) +{ +} template void TrackerTraitsGPU::setBz(float bz) @@ -1121,8 +1184,87 @@ void TrackerTraitsGPU::findCellsNeighboursHybrid(const int iteration) template void TrackerTraitsGPU::findRoadsHybrid(const int iteration) { - TrackerTraits::findRoads(iteration); - mTimeFrameGPU->loadRoadsDevice(); + for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { + const int minimumLayer{startLevel - 1}; + std::vector trackSeeds; + for (int startLayer{mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) { + std::vector lastCellId, updatedCellId; + std::vector lastCellSeed, updatedCellSeed; + + processNeighbours(startLayer, startLevel, mTimeFrame->getCells()[startLayer], lastCellId, updatedCellSeed, updatedCellId); + + int level = startLevel; + for (int iLayer{startLayer - 1}; iLayer > 0 && level > 2; --iLayer) { + lastCellSeed.swap(updatedCellSeed); + lastCellId.swap(updatedCellId); + updatedCellSeed.clear(); + updatedCellId.clear(); + processNeighbours(iLayer, --level, lastCellSeed, lastCellId, updatedCellSeed, updatedCellId); + } + trackSeeds.insert(trackSeeds.end(), updatedCellSeed.begin(), updatedCellSeed.end()); + } + mTimeFrameGPU->createTrackITSExtDevice(trackSeeds); + mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds); + + gpu::fitTrackSeedsKernel<<<20, 512>>>( + mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds, + mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo, + mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks, + trackSeeds.size(), // const size_t nSeeds, + mBz, // const float Bz, + startLevel, // const int startLevel, + mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment, + mTrkParams[0].MaxChi2NDF, // float maxChi2NDF, + mTimeFrameGPU->getDevicePropagator()); // const o2::base::Propagator* propagator + + checkGPUError(cudaHostUnregister(trackSeeds.data())); + mTimeFrameGPU->downloadTrackITSExtDevice(); + + auto& tracks = mTimeFrameGPU->getTrackITSExt(); + std::sort(tracks.begin(), tracks.end(), [](const TrackITSExt& a, const TrackITSExt& b) { + return a.getChi2() < b.getChi2(); + }); + + for (auto& track : tracks) { + int nShared = 0; + bool isFirstShared{false}; + for (int iLayer{0}; iLayer < mTrkParams[0].NLayers; ++iLayer) { + if (track.getClusterIndex(iLayer) == constants::its::UnusedIndex) { + continue; + } + nShared += int(mTimeFrame->isClusterUsed(iLayer, track.getClusterIndex(iLayer))); + isFirstShared |= !iLayer && mTimeFrame->isClusterUsed(iLayer, track.getClusterIndex(iLayer)); + } + + if (nShared > mTrkParams[0].ClusterSharing) { + continue; + } + + std::array rofs{INT_MAX, INT_MAX, INT_MAX}; + for (int iLayer{0}; iLayer < mTrkParams[0].NLayers; ++iLayer) { + if (track.getClusterIndex(iLayer) == constants::its::UnusedIndex) { + continue; + } + mTimeFrame->markUsedCluster(iLayer, track.getClusterIndex(iLayer)); + int currentROF = mTimeFrame->getClusterROF(iLayer, track.getClusterIndex(iLayer)); + for (int iR{0}; iR < 3; ++iR) { + if (rofs[iR] == INT_MAX) { + rofs[iR] = currentROF; + } + if (rofs[iR] == currentROF) { + break; + } + } + } + if (rofs[2] != INT_MAX) { + continue; + } + if (rofs[1] != INT_MAX) { + track.setNextROFbit(); + } + mTimeFrame->getTracks(std::min(rofs[0], rofs[1])).emplace_back(track); + } + } }; template diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h index b009f5fd7ad0a..f2995623adfa4 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h @@ -78,7 +78,8 @@ GPUdi() Cell::Cell(const int firstClusterIndex, const int secondClusterIndex, co class CellSeed final : public o2::track::TrackParCovF { public: - GPUhd() CellSeed() = default; + GPUhdDefault() CellSeed() = default; + GPUhdDefault() CellSeed(const CellSeed&) = default; GPUd() CellSeed(int innerL, int cl0, int cl1, int cl2, int trkl0, int trkl1, o2::track::TrackParCovF& tpc, float chi2) : o2::track::TrackParCovF{tpc}, mChi2{chi2}, mLevel{1} { setUserField(innerL); diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index cf1d13faeaa15..e4b30cf225789 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -132,9 +132,9 @@ void Tracker::clustersToTracksHybrid(std::function logger, } total += evaluateTask(&Tracker::findCellsNeighboursHybrid, "Hybrid Neighbour finding", logger, iteration); logger(fmt::format("\t- Number of Neighbours: {}", mTimeFrame->getNumberOfNeighbours())); - total += evaluateTask(&Tracker::findRoadsHybrid, "Hybrid Road finding", logger, iteration); - logger(fmt::format("\t- Number of Roads: {}", mTimeFrame->getRoads().size())); - total += evaluateTask(&Tracker::findTracksHybrid, "Hybrid Track fitting", logger, iteration); + total += evaluateTask(&Tracker::findRoadsHybrid, "Hybrid Track finding", logger, iteration); + logger(fmt::format("\t- Number of Tracks: {}", mTimeFrame->getNumberOfTracks())); + // total += evaluateTask(&Tracker::findTracksHybrid, "Hybrid Track fitting", logger, iteration); } }