diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 69d6799686654..a4e4328b3aa22 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -209,13 +209,17 @@ void processNeighboursHandler(const int startLayer, template void trackSeedHandler(CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, + const Cluster** unsortedClusters, o2::its::TrackITSExt* tracks, - std::vector& minPtsHost, + const std::vector& layerRadiiHost, + const std::vector& minPtsHost, const unsigned int nSeeds, const float Bz, const int startLevel, - float maxChi2ClusterAttachment, - float maxChi2NDF, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const int reseedIfShorter, + const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, const int nBlocks, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 05810f0074811..f94147747a475 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -325,17 +325,21 @@ void TrackerTraitsGPU::findRoads(const int iteration) mTimeFrameGPU->createTrackITSExtDevice(trackSeeds); mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds); - trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds - mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo - mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks - this->mTrkParams[iteration].MinPt, // std::vector& minPtsHost, + trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* + mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** + mTimeFrameGPU->getDeviceArrayUnsortedClusters(), // Cluster** + mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* + this->mTrkParams[iteration].LayerRadii, // const std::vector& + this->mTrkParams[iteration].MinPt, // const std::vector& trackSeeds.size(), // const size_t nSeeds this->mBz, // const float Bz startLevel, // const int startLevel, this->mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment this->mTrkParams[0].MaxChi2NDF, // float maxChi2NDF - mTimeFrameGPU->getDevicePropagator(), // const o2::base::Propagator* propagator - this->mTrkParams[0].CorrType, // o2::base::PropagatorImpl::MatCorrType + this->mTrkParams[0].ReseedIfShorter, + this->mTrkParams[0].ShiftRefToCluster, + mTimeFrameGPU->getDevicePropagator(), // const o2::base::Propagator* propagator + this->mTrkParams[0].CorrType, // o2::base::PropagatorImpl::MatCorrType conf.nBlocksTracksSeeds[iteration], conf.nThreadsTracksSeeds[iteration]); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 71f1281401e9d..d9136cb96d00e 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -57,35 +57,56 @@ GPUdii() bool fitTrack(TrackITSExt& track, float bz, const TrackingFrameInfo** tfInfos, const o2::base::Propagator* prop, - o2::base::PropagatorF::MatCorrType matCorrType) + o2::base::PropagatorF::MatCorrType matCorrType, + o2::track::TrackPar* linRef, + const bool shiftRefToCluster) { for (int iLayer{start}; iLayer != end; iLayer += step) { if (track.getClusterIndex(iLayer) == constants::UnusedIndex) { continue; } const TrackingFrameInfo& trackingHit = tfInfos[iLayer][track.getClusterIndex(iLayer)]; - if (!track.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame)) { - return false; - } - - if (!prop->propagateToX(track, - trackingHit.xTrackingFrame, - bz, - o2::base::PropagatorImpl::MAX_SIN_PHI, - o2::base::PropagatorImpl::MAX_STEP, - matCorrType)) { - return false; - } + if (linRef) { + if (!track.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame, *linRef, bz)) { + return false; + } + if (!prop->propagateToX(track, + *linRef, + trackingHit.xTrackingFrame, + bz, + o2::base::PropagatorImpl::MAX_SIN_PHI, + o2::base::PropagatorImpl::MAX_STEP, + matCorrType)) { - if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE) { - const float xx0 = (iLayer > 2) ? 1.e-2f : 5.e-3f; // Rough layer thickness - if (!track.correctForMaterial(xx0, xx0 * constants::Radl * constants::Rho, true)) { return false; } + if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE) { + const float xx0 = (iLayer > 2) ? 1.e-2f : 5.e-3f; // Rough layer thickness + if (!track.correctForMaterial(*linRef, xx0, xx0 * constants::Radl * constants::Rho, true)) { + return false; + } + } + } else { + if (!track.o2::track::TrackParCovF::rotate(trackingHit.alphaTrackingFrame)) { + return false; + } + if (!prop->propagateToX(track, + trackingHit.xTrackingFrame, + bz, + o2::base::PropagatorImpl::MAX_SIN_PHI, + o2::base::PropagatorImpl::MAX_STEP, + matCorrType)) { + return false; + } + if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE) { + const float xx0 = (iLayer > 2) ? 1.e-2f : 5.e-3f; // Rough layer thickness + if (!track.correctForMaterial(xx0, xx0 * constants::Radl * constants::Rho, true)) { + return false; + } + } } auto predChi2{track.getPredictedChi2(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)}; - if ((nCl >= 3 && predChi2 > chi2clcut) || predChi2 < 0.f) { return false; } @@ -93,6 +114,10 @@ GPUdii() bool fitTrack(TrackITSExt& track, if (!track.o2::track::TrackParCov::update(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)) { return false; } + if (linRef && shiftRefToCluster) { // displace the reference to the last updated cluster + linRef->setY(trackingHit.positionTrackingFrame[0]); + linRef->setZ(trackingHit.positionTrackingFrame[1]); + } nCl++; } return o2::gpu::CAMath::Abs(track.getQ2Pt()) < maxQoverPt && track.getChi2() < chi2ndfcut * (nCl * 2 - 5); @@ -101,35 +126,85 @@ GPUdii() bool fitTrack(TrackITSExt& track, GPUdii() o2::track::TrackParCov buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3, - const float bz) + const float bz, + const bool reverse = false) { - const float ca = o2::gpu::CAMath::Cos(tf3.alphaTrackingFrame), sa = o2::gpu::CAMath::Sin(tf3.alphaTrackingFrame); + const float sign = reverse ? -1.f : 1.f; + + float ca, sa; + o2::gpu::CAMath::SinCos(tf3.alphaTrackingFrame, sa, ca); + const float x1 = cluster1.xCoordinate * ca + cluster1.yCoordinate * sa; const float y1 = -cluster1.xCoordinate * sa + cluster1.yCoordinate * ca; - const float z1 = cluster1.zCoordinate; const float x2 = cluster2.xCoordinate * ca + cluster2.yCoordinate * sa; const float y2 = -cluster2.xCoordinate * sa + cluster2.yCoordinate * ca; - const float z2 = cluster2.zCoordinate; const float x3 = tf3.xTrackingFrame; const float y3 = tf3.positionTrackingFrame[0]; - const float z3 = tf3.positionTrackingFrame[1]; - - const bool zeroField{o2::gpu::CAMath::Abs(bz) < o2::constants::math::Almost0}; - const float tgp = zeroField ? o2::gpu::CAMath::ATan2(y3 - y1, x3 - x1) : 1.f; - const float crv = zeroField ? 1.f : math_utils::computeCurvature(x3, y3, x2, y2, x1, y1); - const float snp = zeroField ? tgp / o2::gpu::CAMath::Sqrt(1.f + tgp * tgp) : crv * (x3 - math_utils::computeCurvatureCentreX(x3, y3, x2, y2, x1, y1)); - const float tgl12 = math_utils::computeTanDipAngle(x1, y1, x2, y2, z1, z2); - const float tgl23 = math_utils::computeTanDipAngle(x2, y2, x3, y3, z2, z3); - const float q2pt = zeroField ? 1.f / o2::track::kMostProbablePt : crv / (bz * o2::constants::math::B2C); - const float q2pt2 = crv * crv; - const float sg2q2pt = o2::track::kC1Pt2max * (q2pt2 > 0.0005 ? (q2pt2 < 1 ? q2pt2 : 1) : 0.0005); - return track::TrackParCov(tf3.xTrackingFrame, tf3.alphaTrackingFrame, - {y3, z3, snp, 0.5f * (tgl12 + tgl23), q2pt}, - {tf3.covarianceTrackingFrame[0], - tf3.covarianceTrackingFrame[1], tf3.covarianceTrackingFrame[2], - 0.f, 0.f, track::kCSnp2max, - 0.f, 0.f, 0.f, track::kCTgl2max, - 0.f, 0.f, 0.f, 0.f, sg2q2pt}); + + float snp, q2pt, q2pt2; + if (o2::gpu::CAMath::Abs(bz) < 0.01f) { + const float tgp = o2::gpu::CAMath::ATan2(y3 - y1, x3 - x1); + snp = sign * tgp / o2::gpu::CAMath::Sqrt(1.f + tgp * tgp); + q2pt = sign / track::kMostProbablePt; + q2pt2 = 1.f; + } else { + const float crv = math_utils::computeCurvature(x3, y3, x2, y2, x1, y1); + snp = sign * crv * (x3 - math_utils::computeCurvatureCentreX(x3, y3, x2, y2, x1, y1)); + q2pt = sign * crv / (bz * o2::constants::math::B2C); + q2pt2 = crv * crv; + } + + const float tgl = 0.5f * (math_utils::computeTanDipAngle(x1, y1, x2, y2, cluster1.zCoordinate, cluster2.zCoordinate) + + math_utils::computeTanDipAngle(x2, y2, x3, y3, cluster2.zCoordinate, tf3.positionTrackingFrame[1])); + const float sg2q2pt = track::kC1Pt2max * (q2pt2 > 0.0005f ? (q2pt2 < 1.f ? q2pt2 : 1.f) : 0.0005f); + + return {x3, tf3.alphaTrackingFrame, {y3, tf3.positionTrackingFrame[1], snp, tgl, q2pt}, {tf3.covarianceTrackingFrame[0], tf3.covarianceTrackingFrame[1], tf3.covarianceTrackingFrame[2], 0.f, 0.f, track::kCSnp2max, 0.f, 0.f, 0.f, track::kCTgl2max, 0.f, 0.f, 0.f, 0.f, sg2q2pt}}; +} + +template +GPUdii() TrackITSExt seedTrackForRefit(const CellSeed& seed, + const TrackingFrameInfo** foundTrackingFrameInfo, + const Cluster** unsortedClusters, + const float* layerRadii, + const float bz, + const int reseedIfShorter) +{ + TrackITSExt temporaryTrack(seed); + int lrMin = nLayers, lrMax = 0, lrMid = 0; + for (int iL{0}; iL < nLayers; ++iL) { + const int idx = seed.getCluster(iL); + temporaryTrack.setExternalClusterIndex(iL, idx, idx != constants::UnusedIndex); + if (idx != constants::UnusedIndex) { + // TODO only works if does not have holes + lrMin = o2::gpu::CAMath::Min(lrMin, iL); + lrMax = o2::gpu::CAMath::Max(lrMax, iL); + } + } + const int ncl = temporaryTrack.getNClusters(); + if (ncl < reseedIfShorter && ncl > 0) { // need to check if there are any clusters since we keep invalidate seeeds around + if (ncl == nLayers) { + lrMin = 0; + lrMax = nLayers - 1; + lrMid = (lrMin + lrMax) / 2; + } else { + lrMid = lrMin + 1; + float midR = 0.5f * (layerRadii[lrMax] + layerRadii[lrMin]), dstMidR = o2::gpu::CAMath::Abs(midR - layerRadii[lrMid]); + for (int iL = lrMid + 1; iL < lrMax; ++iL) { // find the midpoint as closest to the midR + auto dst = o2::gpu::GPUCommonMath::Abs(midR - layerRadii[iL]); + if (dst < dstMidR) { + lrMid = iL; + dstMidR = dst; + } + } + } + const auto& cluster0_tf = foundTrackingFrameInfo[lrMin][seed.getCluster(lrMin)]; + const auto& cluster1_gl = unsortedClusters[lrMid][seed.getCluster(lrMid)]; + const auto& cluster2_gl = unsortedClusters[lrMax][seed.getCluster(lrMax)]; + temporaryTrack.getParamIn() = buildTrackSeed(cluster2_gl, cluster1_gl, cluster0_tf, bz, true); + } + temporaryTrack.resetCovariance(); + temporaryTrack.setCov(temporaryTrack.getQ2Pt() * temporaryTrack.getQ2Pt() * temporaryTrack.getCov()[o2::track::CovLabels::kSigQ2Pt2], o2::track::CovLabels::kSigQ2Pt2); + return temporaryTrack; } struct sort_tracklets { @@ -206,27 +281,23 @@ template GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, + const Cluster** unsortedClusters, o2::its::TrackITSExt* tracks, + const float* layerRadii, const float* minPts, const unsigned int nSeeds, const float bz, const int startLevel, - float maxChi2ClusterAttachment, - float maxChi2NDF, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const int reseedIfShorter, + const bool shifRefToCluster, 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) { - auto& seed = trackSeeds[iCurrentTrackSeedIndex]; - - TrackITSExt temporaryTrack{seed}; - - temporaryTrack.resetCovariance(); - temporaryTrack.setChi2(0); - auto& clusters = seed.getClusters(); - for (int iL{0}; iL < nLayers; ++iL) { - temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::UnusedIndex); - } + TrackITSExt temporaryTrack = seedTrackForRefit(trackSeeds[iCurrentTrackSeedIndex], foundTrackingFrameInfo, unsortedClusters, layerRadii, bz, reseedIfShorter); + o2::track::TrackPar linRef{temporaryTrack}; bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, 0, // int lastLayer, nLayers, // int firstLayer, @@ -238,14 +309,17 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( bz, // float bz, foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, propagator, // const o2::base::Propagator* propagator, - matCorrType); // o2::base::PropagatorF::MatCorrType matCorrType + matCorrType, // o2::base::PropagatorF::MatCorrType matCorrType + &linRef, + shifRefToCluster); if (!fitSuccess) { continue; } temporaryTrack.getParamOut() = temporaryTrack.getParamIn(); + linRef = temporaryTrack.getParamOut(); // use refitted track as lin.reference temporaryTrack.resetCovariance(); + temporaryTrack.setCov(temporaryTrack.getQ2Pt() * temporaryTrack.getQ2Pt() * temporaryTrack.getCov()[o2::track::CovLabels::kSigQ2Pt2], o2::track::CovLabels::kSigQ2Pt2); temporaryTrack.setChi2(0); - fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, nLayers - 1, // int lastLayer, -1, // int firstLayer, @@ -257,7 +331,9 @@ GPUg() void __launch_bounds__(256, 1) fitTrackSeedsKernel( bz, // float bz, foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, propagator, // const o2::base::Propagator* propagator, - matCorrType); // o2::base::PropagatorF::MatCorrType matCorrType + matCorrType, // o2::base::PropagatorF::MatCorrType matCorrType + &linRef, + shifRefToCluster); if (!fitSuccess || temporaryTrack.getPt() < minPts[nLayers - temporaryTrack.getNClusters()]) { continue; } @@ -1088,34 +1164,42 @@ void processNeighboursHandler(const int startLayer, template void trackSeedHandler(CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, + const Cluster** unsortedClusters, o2::its::TrackITSExt* tracks, - std::vector& minPtsHost, + const std::vector& layerRadiiHost, + const std::vector& minPtsHost, const unsigned int nSeeds, const float bz, const int startLevel, - float maxChi2ClusterAttachment, - float maxChi2NDF, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const int reseedIfShorter, + const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, const int nBlocks, const int nThreads) { thrust::device_vector minPts(minPtsHost); + thrust::device_vector layerRadii(layerRadiiHost); gpu::fitTrackSeedsKernel<<>>( - trackSeeds, // CellSeed* - foundTrackingFrameInfo, // TrackingFrameInfo** - tracks, // TrackITSExt* - thrust::raw_pointer_cast(&minPts[0]), // const float* minPts, - nSeeds, // const unsigned int - bz, // const float - startLevel, // const int - maxChi2ClusterAttachment, // float - maxChi2NDF, // float - propagator, // const o2::base::Propagator* - matCorrType); // o2::base::PropagatorF::MatCorrType + trackSeeds, // CellSeed* + foundTrackingFrameInfo, // TrackingFrameInfo** + unsortedClusters, // Cluster** + tracks, // TrackITSExt* + thrust::raw_pointer_cast(&layerRadii[0]), // const float* + thrust::raw_pointer_cast(&minPts[0]), // const float* + nSeeds, // const unsigned int + bz, // const float + startLevel, // const int + maxChi2ClusterAttachment, // float + maxChi2NDF, // float + reseedIfShorter, // int + shiftRefToCluster, // bool + propagator, // const o2::base::Propagator* + matCorrType); // o2::base::PropagatorF::MatCorrType thrust::device_ptr tr_ptr(tracks); thrust::sort(tr_ptr, tr_ptr + nSeeds, gpu::compare_track_chi2()); - GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream)); } /// Explicit instantiation of ITS2 handlers @@ -1281,13 +1365,17 @@ template void processNeighboursHandler<7>(const int startLayer, template void trackSeedHandler(CellSeed<7>* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, + const Cluster** unsortedClusters, o2::its::TrackITSExt* tracks, - std::vector& minPtsHost, + const std::vector& layerRadiiHost, + const std::vector& minPtsHost, const unsigned int nSeeds, const float bz, const int startLevel, - float maxChi2ClusterAttachment, - float maxChi2NDF, + const float maxChi2ClusterAttachment, + const float maxChi2NDF, + const int reseedIfShorter, + const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, const int nBlocks, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 9a6452270d144..000c8fe822498 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -66,10 +66,10 @@ struct TrackingParameters { o2::base::PropagatorImpl::MatCorrType CorrType = o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE; float MaxChi2ClusterAttachment = 60.f; float MaxChi2NDF = 30.f; - int reseedIfShorter = 6; // reseed for the final fit track with the length shorter than this + int ReseedIfShorter = 6; // reseed for the final fit track with the length shorter than this std::vector MinPt = {0.f, 0.f, 0.f, 0.f}; unsigned char StartLayerMask = 0x7F; - bool shiftRefToCluster = true; // TrackFit: after update shift the linearization reference to cluster + bool ShiftRefToCluster = true; // TrackFit: after update shift the linearization reference to cluster bool FindShortTracks = false; bool PerPrimaryVertexProcessing = false; bool SaveTimeBenchmarks = false; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index f582b5ef3aec5..ddc32ed18cbfe 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -92,7 +92,7 @@ class TrackerTraits virtual int getTFNumberOfCells() const { return mTimeFrame->getNumberOfCells(); } private: - track::TrackParCov buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3); + track::TrackParCov buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3, bool reverse = false); TrackITSExt seedTrackForRefit(const CellSeedN& seed); bool fitTrack(TrackITSExt& track, int start, int end, int step, float chi2clcut = o2::constants::math::VeryBig, float chi2ndfcut = o2::constants::math::VeryBig, float maxQoverPt = o2::constants::math::VeryBig, int nCl = 0, o2::track::TrackPar* refLin = nullptr); diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx index c6ed343033996..87787eeee03a9 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx @@ -186,8 +186,8 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode int lslot = tc.MaxTrackLength - ilg; p.MinPt[lslot] *= bFactor; } - p.reseedIfShorter = tc.reseedIfShorter; - p.shiftRefToCluster = tc.shiftRefToCluster; + p.ReseedIfShorter = tc.reseedIfShorter; + p.ShiftRefToCluster = tc.shiftRefToCluster; p.createArtefactLabels = tc.createArtefactLabels; p.PrintMemory = tc.printMemory; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 5c5eb47216051..6b237ad0a63e8 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -765,7 +765,6 @@ void TrackerTraits::findRoads(const int iteration) auto forSeed = [&](auto Tag, int iSeed, int offset = 0) { TrackITSExt temporaryTrack = seedTrackForRefit(trackSeeds[iSeed]); o2::track::TrackPar linRef{temporaryTrack}; - o2::track::TrackParCov savTr = temporaryTrack; // REMOVE bool fitSuccess = fitTrack(temporaryTrack, 0, mTrkParams[0].NLayers, 1, mTrkParams[0].MaxChi2ClusterAttachment, mTrkParams[0].MaxChi2NDF, o2::constants::math::VeryBig, 0, &linRef); if (!fitSuccess) { return 0; @@ -773,7 +772,7 @@ void TrackerTraits::findRoads(const int iteration) temporaryTrack.getParamOut() = temporaryTrack.getParamIn(); linRef = temporaryTrack.getParamOut(); // use refitted track as lin.reference temporaryTrack.resetCovariance(); - temporaryTrack.setCov(temporaryTrack.getQ2Pt() * temporaryTrack.getQ2Pt() * temporaryTrack.getCov()[14], 14); + temporaryTrack.setCov(temporaryTrack.getQ2Pt() * temporaryTrack.getQ2Pt() * temporaryTrack.getCov()[o2::track::CovLabels::kSigQ2Pt2], o2::track::CovLabels::kSigQ2Pt2); temporaryTrack.setChi2(0); fitSuccess = fitTrack(temporaryTrack, mTrkParams[0].NLayers - 1, -1, -1, mTrkParams[0].MaxChi2ClusterAttachment, mTrkParams[0].MaxChi2NDF, 50.f, 0, &linRef); if (!fitSuccess || temporaryTrack.getPt() < mTrkParams[iteration].MinPt[mTrkParams[iteration].NLayers - temporaryTrack.getNClusters()]) { @@ -1082,7 +1081,7 @@ bool TrackerTraits::fitTrack(TrackITSExt& track, int start, int end, in if (!track.o2::track::TrackParCov::update(trackingHit.positionTrackingFrame, trackingHit.covarianceTrackingFrame)) { return false; } - if (linRef && mTrkParams[0].shiftRefToCluster) { // displace the reference to the last updated cluster + if (linRef && mTrkParams[0].ShiftRefToCluster) { // displace the reference to the last updated cluster linRef->setY(trackingHit.positionTrackingFrame[0]); linRef->setZ(trackingHit.positionTrackingFrame[1]); } @@ -1208,27 +1207,22 @@ template TrackITSExt TrackerTraits::seedTrackForRefit(const CellSeedN& seed) { TrackITSExt temporaryTrack(seed); + int lrMin = nLayers, lrMax = 0, lrMid = 0; for (int iL = 0; iL < nLayers; ++iL) { - temporaryTrack.setExternalClusterIndex(iL, seed.getCluster(iL), seed.getCluster(iL) != constants::UnusedIndex); + const int idx = seed.getCluster(iL); + temporaryTrack.setExternalClusterIndex(iL, idx, idx != constants::UnusedIndex); + if (idx != constants::UnusedIndex) { + lrMin = o2::gpu::CAMath::Min(lrMin, iL); + lrMax = o2::gpu::CAMath::Max(lrMax, iL); + } } int ncl = temporaryTrack.getNClusters(); - if (ncl < mTrkParams[0].reseedIfShorter) { // reseed with circle passing via edges and the midpoint - int lrMin = 999, lrMax = 0, lrMid = 0; + if (ncl < mTrkParams[0].ReseedIfShorter) { // reseed with circle passing via edges and the midpoint if (ncl == mTrkParams[0].NLayers) { lrMin = 0; lrMax = mTrkParams[0].NLayers - 1; lrMid = (lrMin + lrMax) / 2; } else { - for (int iL = 0; iL < nLayers; ++iL) { - if (seed.getCluster(iL) != constants::UnusedIndex) { - if (iL < lrMin) { - lrMin = iL; - } - if (iL > lrMax) { - lrMax = iL; - } - } - } lrMid = lrMin + 1; float midR = 0.5 * (mTrkParams[0].LayerRadii[lrMax] + mTrkParams[0].LayerRadii[lrMin]), dstMidR = o2::gpu::GPUCommonMath::Abs(midR - mTrkParams[0].LayerRadii[lrMid]); for (int iL = lrMid + 1; iL < lrMax; ++iL) { // find the midpoint as closest to the midR @@ -1242,45 +1236,48 @@ TrackITSExt TrackerTraits::seedTrackForRefit(const CellSeedN& seed) const auto& cluster0_tf = mTimeFrame->getTrackingFrameInfoOnLayer(lrMin)[seed.getCluster(lrMin)]; // if the sensor frame! const auto& cluster1_gl = mTimeFrame->getUnsortedClusters()[lrMid][seed.getCluster(lrMid)]; // global frame const auto& cluster2_gl = mTimeFrame->getUnsortedClusters()[lrMax][seed.getCluster(lrMax)]; // global frame - temporaryTrack.getParamIn() = buildTrackSeed(cluster2_gl, cluster1_gl, cluster0_tf); - temporaryTrack.setQ2Pt(-temporaryTrack.getQ2Pt()); // we are calling buildTrackSeed with the clusters order opposite to what it expects - temporaryTrack.setSnp(-temporaryTrack.getSnp()); // we are calling buildTrackSeed with the clusters order opposite to what it expects + temporaryTrack.getParamIn() = buildTrackSeed(cluster2_gl, cluster1_gl, cluster0_tf, true); } temporaryTrack.resetCovariance(); - temporaryTrack.setCov(temporaryTrack.getQ2Pt() * temporaryTrack.getQ2Pt() * temporaryTrack.getCov()[14], 14); + temporaryTrack.setCov(temporaryTrack.getQ2Pt() * temporaryTrack.getQ2Pt() * temporaryTrack.getCov()[o2::track::CovLabels::kSigQ2Pt2], o2::track::CovLabels::kSigQ2Pt2); return temporaryTrack; } /// Clusters are given from inside outward (cluster3 is the outermost). The outermost cluster is given in the tracking /// frame coordinates whereas the others are referred to the global frame. template -track::TrackParCov TrackerTraits::buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3) +track::TrackParCov TrackerTraits::buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3, bool reverse) { - float ca{-999.f}, sa{-999.f}; + const float sign = reverse ? -1.f : 1.f; + + float ca, sa; o2::gpu::CAMath::SinCos(tf3.alphaTrackingFrame, sa, ca); + const float x1 = cluster1.xCoordinate * ca + cluster1.yCoordinate * sa; const float y1 = -cluster1.xCoordinate * sa + cluster1.yCoordinate * ca; - const float z1 = cluster1.zCoordinate; const float x2 = cluster2.xCoordinate * ca + cluster2.yCoordinate * sa; const float y2 = -cluster2.xCoordinate * sa + cluster2.yCoordinate * ca; - const float z2 = cluster2.zCoordinate; const float x3 = tf3.xTrackingFrame; const float y3 = tf3.positionTrackingFrame[0]; - const float z3 = tf3.positionTrackingFrame[1]; - float tgp{1.f}, crv{1.f}, snp{-999.f}, tgl12{-999.f}, tgl23{-999.f}, q2pt{1.f / track::kMostProbablePt}, q2pt2{1.f}, sg2q2pt{-999.f}; + + float snp, q2pt, q2pt2; if (mIsZeroField) { - tgp = o2::gpu::CAMath::ATan2(y3 - y1, x3 - x1); - snp = tgp / o2::gpu::CAMath::Sqrt(1.f + tgp * tgp); + const float tgp = o2::gpu::CAMath::ATan2(y3 - y1, x3 - x1); + snp = sign * tgp / o2::gpu::CAMath::Sqrt(1.f + tgp * tgp); + q2pt = sign / track::kMostProbablePt; + q2pt2 = 1.f; } else { - crv = math_utils::computeCurvature(x3, y3, x2, y2, x1, y1); - snp = crv * (x3 - math_utils::computeCurvatureCentreX(x3, y3, x2, y2, x1, y1)); - q2pt = crv / (mBz * o2::constants::math::B2C); + const float crv = math_utils::computeCurvature(x3, y3, x2, y2, x1, y1); + snp = sign * crv * (x3 - math_utils::computeCurvatureCentreX(x3, y3, x2, y2, x1, y1)); + q2pt = sign * crv / (mBz * o2::constants::math::B2C); q2pt2 = crv * crv; } - tgl12 = math_utils::computeTanDipAngle(x1, y1, x2, y2, z1, z2); - tgl23 = math_utils::computeTanDipAngle(x2, y2, x3, y3, z2, z3); - sg2q2pt = track::kC1Pt2max * (q2pt2 > 0.0005f ? (q2pt2 < 1.f ? q2pt2 : 1.f) : 0.0005f); - return {tf3.xTrackingFrame, tf3.alphaTrackingFrame, {y3, z3, snp, 0.5f * (tgl12 + tgl23), q2pt}, {tf3.covarianceTrackingFrame[0], tf3.covarianceTrackingFrame[1], tf3.covarianceTrackingFrame[2], 0.f, 0.f, track::kCSnp2max, 0.f, 0.f, 0.f, track::kCTgl2max, 0.f, 0.f, 0.f, 0.f, sg2q2pt}}; + + const float tgl = 0.5f * (math_utils::computeTanDipAngle(x1, y1, x2, y2, cluster1.zCoordinate, cluster2.zCoordinate) + + math_utils::computeTanDipAngle(x2, y2, x3, y3, cluster2.zCoordinate, tf3.positionTrackingFrame[1])); + const float sg2q2pt = track::kC1Pt2max * (q2pt2 > 0.0005f ? (q2pt2 < 1.f ? q2pt2 : 1.f) : 0.0005f); + + return {x3, tf3.alphaTrackingFrame, {y3, tf3.positionTrackingFrame[1], snp, tgl, q2pt}, {tf3.covarianceTrackingFrame[0], tf3.covarianceTrackingFrame[1], tf3.covarianceTrackingFrame[2], 0.f, 0.f, track::kCSnp2max, 0.f, 0.f, 0.f, track::kCTgl2max, 0.f, 0.f, 0.f, 0.f, sg2q2pt}}; } template