diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index a058f7e5fab0c..b0fb443513fef 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -131,6 +131,7 @@ void countCellsHandler(const Cluster** sortedClusters, CellSeed* cells, int** cellsLUTsDeviceArray, int* cellsLUTsHost, + const int deltaROF, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, @@ -148,6 +149,7 @@ void computeCellsHandler(const Cluster** sortedClusters, CellSeed* cells, int** cellsLUTsDeviceArray, int* cellsLUTsHost, + const int deltaROF, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, @@ -160,6 +162,8 @@ unsigned int countCellNeighboursHandler(CellSeed** cellsLayersDevice, int** cellsLUTs, gpuPair* cellNeighbours, int* neighboursIndexTable, + const Tracklet** tracklets, + const int deltaROF, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, @@ -174,6 +178,8 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice, int** cellsLUTs, gpuPair* cellNeighbours, int* neighboursIndexTable, + const Tracklet** tracklets, + const int deltaROF, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index a8061e872c029..46ae8f49d6145 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -140,6 +140,7 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) nullptr, mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceCellLUTs()[iLayer], + this->mTrkParams[iteration].DeltaROF, this->mBz, this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].CellDeltaTanLambdaSigma, @@ -157,6 +158,7 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) mTimeFrameGPU->getDeviceCells()[iLayer], mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceCellLUTs()[iLayer], + this->mTrkParams[iteration].DeltaROF, this->mBz, this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].CellDeltaTanLambdaSigma, @@ -185,6 +187,8 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceNeighbourPairs(iLayer), mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer), + (const Tracklet**)mTimeFrameGPU->getDeviceArrayTracklets(), + this->mTrkParams[0].DeltaROF, this->mTrkParams[0].MaxChi2ClusterAttachment, this->mBz, iLayer, @@ -201,6 +205,8 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceNeighbourPairs(iLayer), mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer), + (const Tracklet**)mTimeFrameGPU->getDeviceArrayTracklets(), + this->mTrkParams[0].DeltaROF, this->mTrkParams[0].MaxChi2ClusterAttachment, this->mBz, iLayer, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index fb75764da2e36..02be19b1e3a08 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -201,7 +201,13 @@ GPUd() o2::track::TrackParCov buildTrackSeed(const Cluster& cluster1, } struct sort_tracklets { - GPUhd() bool operator()(const Tracklet& a, const Tracklet& b) { return a.firstClusterIndex < b.firstClusterIndex || (a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex < b.secondClusterIndex); } + GPUhd() bool operator()(const Tracklet& a, const Tracklet& b) + { + if (a.firstClusterIndex != b.firstClusterIndex) { + return a.firstClusterIndex < b.firstClusterIndex; + } + return a.secondClusterIndex < b.secondClusterIndex; + } }; struct equal_tracklets { @@ -263,23 +269,34 @@ struct compare_track_chi2 { } }; -GPUd() gpuSpan getPrimaryVertices(const int rof, - const int* roframesPV, - const int nROF, - const uint8_t* mask, - const Vertex* vertices) +GPUdii() gpuSpan getPrimaryVertices(const int rof, + const int* roframesPV, + const int nROF, + const uint8_t* mask, + const Vertex* vertices) { const int start_pv_id = roframesPV[rof]; const int stop_rof = rof >= nROF - 1 ? nROF : rof + 1; - size_t delta = mask[rof] ? roframesPV[stop_rof] - start_pv_id : 0; // return empty span if ROF is excluded + const size_t delta = mask[rof] ? roframesPV[stop_rof] - start_pv_id : 0; // return empty span if ROF is excluded return gpuSpan(&vertices[start_pv_id], delta); }; -GPUd() gpuSpan getClustersOnLayer(const int rof, - const int totROFs, - const int layer, - const int** roframesClus, - const Cluster** clusters) +GPUdii() gpuSpan getPrimaryVertices(const int romin, + const int romax, + const int* roframesPV, + const int nROF, + const Vertex* vertices) +{ + const int start_pv_id = roframesPV[romin]; + const int stop_rof = romax >= nROF - 1 ? nROF : romax + 1; + return gpuSpan(&vertices[start_pv_id], roframesPV[stop_rof] - roframesPV[romin]); +}; + +GPUdii() gpuSpan getClustersOnLayer(const int rof, + const int totROFs, + const int layer, + const int** roframesClus, + const Cluster** clusters) { if (rof < 0 || rof >= totROFs) { return gpuSpan(); @@ -360,6 +377,8 @@ GPUg() void computeLayerCellNeighboursKernel( int* neighboursIndexTable, int** cellsLUTs, gpuPair* cellNeighbours, + const Tracklet** tracklets, + const int deltaROF, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, @@ -377,15 +396,29 @@ GPUg() void computeLayerCellNeighboursKernel( if (nextCellSeed.getFirstTrackletIndex() != nextLayerTrackletIndex) { // Check if cells share the same tracklet break; } + + if (deltaROF) { + const auto& trkl00 = tracklets[layerIndex][currentCellSeed.getFirstTrackletIndex()]; + const auto& trkl01 = tracklets[layerIndex + 1][currentCellSeed.getSecondTrackletIndex()]; + const auto& trkl10 = tracklets[layerIndex + 1][nextCellSeed.getFirstTrackletIndex()]; + const auto& trkl11 = tracklets[layerIndex + 2][nextCellSeed.getSecondTrackletIndex()]; + if ((o2::gpu::CAMath::Max(trkl00.getMaxRof(), o2::gpu::CAMath::Max(trkl01.getMaxRof(), o2::gpu::CAMath::Max(trkl10.getMaxRof(), trkl11.getMaxRof()))) - + o2::gpu::CAMath::Min(trkl00.getMinRof(), o2::gpu::CAMath::Min(trkl01.getMinRof(), o2::gpu::CAMath::Min(trkl10.getMinRof(), trkl11.getMinRof())))) > deltaROF) { + continue; + } + } + if (!nextCellSeed.rotate(currentCellSeed.getAlpha()) || !nextCellSeed.propagateTo(currentCellSeed.getX(), bz)) { continue; } + float chi2 = currentCellSeed.getPredictedChi2(nextCellSeed); if (chi2 > maxChi2ClusterAttachment) /// TODO: switch to the chi2 wrt cluster to avoid correlation { continue; } + if constexpr (initRun) { atomicAdd(neighboursLUT + iNextCell, 1); neighboursIndexTable[iCurrentCellIndex]++; @@ -412,6 +445,7 @@ GPUg() void computeLayerCellsKernel( const int layer, CellSeed* cells, int** cellsLUTs, + const int deltaROF, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, @@ -432,6 +466,9 @@ GPUg() void computeLayerCellsKernel( break; } const Tracklet& nextTracklet = tracklets[layer + 1][iNextTrackletIndex]; + if (deltaROF && currentTracklet.getSpanRof(nextTracklet) > deltaROF) { + continue; + } const float deltaTanLambda{o2::gpu::CAMath::Abs(currentTracklet.tanLambda - nextTracklet.tanLambda)}; if (deltaTanLambda / cellDeltaTanLambdaSigma < nSigmaCut) { @@ -515,9 +552,12 @@ GPUg() void computeLayerTrackletsMultiROFKernel( { const int phiBins{utils->getNphiBins()}; const int zBins{utils->getNzBins()}; + const int tableSize{phiBins * zBins + 1}; for (unsigned int iROF{blockIdx.x}; iROF < endROF - startROF; iROF += gridDim.x) { - const short rof0 = iROF + startROF; - auto primaryVertices = getPrimaryVertices(rof0, rofPV, totalROFs, multMask, vertices); + const short pivotROF = iROF + startROF; + const short minROF = o2::gpu::CAMath::Max(startROF, static_cast(pivotROF - deltaROF)); + const short maxROF = o2::gpu::CAMath::Min(endROF - 1, static_cast(pivotROF + deltaROF)); + auto primaryVertices = getPrimaryVertices(minROF, maxROF, rofPV, totalROFs, vertices); if (primaryVertices.empty()) { continue; } @@ -526,17 +566,17 @@ GPUg() void computeLayerTrackletsMultiROFKernel( if ((endVtx - startVtx) <= 0) { continue; } - const short minROF = o2::gpu::CAMath::Max(startROF, static_cast(rof0 - deltaROF)); - const short maxROF = o2::gpu::CAMath::Min(endROF - 1, static_cast(rof0 + deltaROF)); - auto clustersCurrentLayer = getClustersOnLayer(rof0, totalROFs, layerIndex, ROFClusters, clusters); + + auto clustersCurrentLayer = getClustersOnLayer(pivotROF, totalROFs, layerIndex, ROFClusters, clusters); if (clustersCurrentLayer.empty()) { continue; } for (int currentClusterIndex = threadIdx.x; currentClusterIndex < clustersCurrentLayer.size(); currentClusterIndex += blockDim.x) { + unsigned int storedTracklets{0}; const auto& currentCluster{clustersCurrentLayer[currentClusterIndex]}; - const int currentSortedIndex{ROFClusters[layerIndex][rof0] + currentClusterIndex}; + const int currentSortedIndex{ROFClusters[layerIndex][pivotROF] + currentClusterIndex}; if (usedClusters[layerIndex][currentCluster.clusterId]) { continue; } @@ -564,9 +604,8 @@ GPUg() void computeLayerTrackletsMultiROFKernel( phiBinsNum += phiBins; } - const int tableSize{phiBins * zBins + 1}; - for (short rof1{minROF}; rof1 <= maxROF; ++rof1) { - auto clustersNextLayer = getClustersOnLayer(rof1, totalROFs, layerIndex + 1, ROFClusters, clusters); + for (short targetROF{minROF}; targetROF <= maxROF; ++targetROF) { + auto clustersNextLayer = getClustersOnLayer(targetROF, totalROFs, layerIndex + 1, ROFClusters, clusters); if (clustersNextLayer.empty()) { continue; } @@ -574,8 +613,8 @@ GPUg() void computeLayerTrackletsMultiROFKernel( int iPhiBin = (selectedBinsRect.y + iPhiCount) % phiBins; const int firstBinIndex{utils->getBinIndex(selectedBinsRect.x, iPhiBin)}; const int maxBinIndex{firstBinIndex + selectedBinsRect.z - selectedBinsRect.x + 1}; - const int firstRowClusterIndex = indexTables[layerIndex + 1][(rof1)*tableSize + firstBinIndex]; - const int maxRowClusterIndex = indexTables[layerIndex + 1][(rof1)*tableSize + maxBinIndex]; + const int firstRowClusterIndex = indexTables[layerIndex + 1][(targetROF)*tableSize + firstBinIndex]; + const int maxRowClusterIndex = indexTables[layerIndex + 1][(targetROF)*tableSize + maxBinIndex]; for (int nextClusterIndex{firstRowClusterIndex}; nextClusterIndex < maxRowClusterIndex; ++nextClusterIndex) { if (nextClusterIndex >= clustersNextLayer.size()) { break; @@ -592,8 +631,8 @@ GPUg() void computeLayerTrackletsMultiROFKernel( } else { const float phi{o2::gpu::CAMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate)}; const float tanL{(currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius)}; - const int nextSortedIndex{ROFClusters[layerIndex + 1][rof1] + nextClusterIndex}; - new (tracklets[layerIndex] + trackletsLUT[layerIndex][currentSortedIndex] + storedTracklets) Tracklet{currentSortedIndex, nextSortedIndex, tanL, phi, rof0, rof1}; + const int nextSortedIndex{ROFClusters[layerIndex + 1][targetROF] + nextClusterIndex}; + new (tracklets[layerIndex] + trackletsLUT[layerIndex][currentSortedIndex] + storedTracklets) Tracklet{currentSortedIndex, nextSortedIndex, tanL, phi, pivotROF, targetROF}; } ++storedTracklets; } @@ -1018,6 +1057,7 @@ void countCellsHandler( CellSeed* cells, int** cellsLUTsArrayDevice, int* cellsLUTsHost, + const int deltaROF, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, @@ -1035,6 +1075,7 @@ void countCellsHandler( layer, // const int cells, // CellSeed* cellsLUTsArrayDevice, // int** + deltaROF, // const int bz, // const float maxChi2ClusterAttachment, // const float cellDeltaTanLambdaSigma, // const float @@ -1053,6 +1094,7 @@ void computeCellsHandler( CellSeed* cells, int** cellsLUTsArrayDevice, int* cellsLUTsHost, + const int deltaROF, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, @@ -1070,6 +1112,7 @@ void computeCellsHandler( layer, // const int cells, // CellSeed* cellsLUTsArrayDevice, // int** + deltaROF, // const int bz, // const float maxChi2ClusterAttachment, // const float cellDeltaTanLambdaSigma, // const float @@ -1081,6 +1124,8 @@ unsigned int countCellNeighboursHandler(CellSeed** cellsLayersDevice, int** cellsLUTs, gpuPair* cellNeighbours, int* neighboursIndexTable, + const Tracklet** tracklets, + const int deltaROF, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, @@ -1096,12 +1141,13 @@ unsigned int countCellNeighboursHandler(CellSeed** cellsLayersDevice, neighboursIndexTable, cellsLUTs, cellNeighbours, + tracklets, + deltaROF, maxChi2ClusterAttachment, bz, layerIndex, nCells, maxCellNeighbours); - gpu::cubInclusiveScanInPlace(neighboursLUT, nCellsNext); gpu::cubExclusiveScanInPlace(neighboursIndexTable, nCells + 1); unsigned int nNeighbours; @@ -1114,6 +1160,8 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice, int** cellsLUTs, gpuPair* cellNeighbours, int* neighboursIndexTable, + const Tracklet** tracklets, + const int deltaROF, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, @@ -1130,6 +1178,8 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice, neighboursIndexTable, cellsLUTs, cellNeighbours, + tracklets, + deltaROF, maxChi2ClusterAttachment, bz, layerIndex, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h index 5741a9fc65947..e6c9db55198a3 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracklet.h @@ -47,13 +47,10 @@ struct Tracklet final { GPUhdi() auto getDeltaRof() const { return rof[1] - rof[0]; } GPUhdi() auto getSpanRof(const Tracklet& o) const noexcept { return o2::gpu::CAMath::Max(getMaxRof(), o.getMaxRof()) - o2::gpu::CAMath::Min(getMinRof(), o.getMinRof()); } GPUhdi() unsigned char operator<(const Tracklet&) const; -#if !defined(GPUCA_NO_FMT) && !defined(GPUCA_GPUCODE_DEVICE) - std::string asString() const + GPUhd() void print() const { - return fmt::format("fClIdx:{} fROF:{} sClIdx:{} sROF:{} (DROF:{})", firstClusterIndex, rof[0], secondClusterIndex, rof[1], getDeltaRof()); + printf("TRKLT: fClIdx:%d fROF:%d sClIdx:%d sROF:%d (DROF:%d) tgl=%f phi=%f\n", firstClusterIndex, rof[0], secondClusterIndex, rof[1], getDeltaRof(), tanLambda, phi); } - void print() const { LOG(info) << asString(); } -#endif int firstClusterIndex{constants::UnusedIndex}; int secondClusterIndex{constants::UnusedIndex}; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index b46e7a68875e6..2515f8287d72b 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -73,6 +73,9 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROF mTaskArena->execute([&] { auto forTracklets = [&](auto Tag, int iLayer, int pivotROF, int base, int& offset) -> int { + if (!mTimeFrame->mMultiplicityCutMask[pivotROF]) { + return 0; + } int minROF = o2::gpu::CAMath::Max(startROF, pivotROF - mTrkParams[iteration].DeltaROF); int maxROF = o2::gpu::CAMath::Min(endROF - 1, pivotROF + mTrkParams[iteration].DeltaROF); gsl::span primaryVertices = mTrkParams[iteration].UseDiamond ? diamondSpan : mTimeFrame->getPrimaryVertices(minROF, maxROF); @@ -87,103 +90,98 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROF int localCount = 0; auto& tracklets = mTimeFrame->getTracklets()[iLayer]; - for (int targetROF0{minROF}; targetROF0 <= maxROF; ++targetROF0) { - if (!mTimeFrame->mMultiplicityCutMask[targetROF0]) { - continue; - } - auto layer0 = mTimeFrame->getClustersOnLayer(targetROF0, iLayer); - if (layer0.empty()) { + auto layer0 = mTimeFrame->getClustersOnLayer(pivotROF, iLayer); + if (layer0.empty()) { + return 0; + } + + const float meanDeltaR = mTrkParams[iteration].LayerRadii[iLayer + 1] - mTrkParams[iteration].LayerRadii[iLayer]; + + for (int iCluster = 0; iCluster < int(layer0.size()); ++iCluster) { + const Cluster& currentCluster = layer0[iCluster]; + const int currentSortedIndex = mTimeFrame->getSortedIndex(pivotROF, iLayer, iCluster); + if (mTimeFrame->isClusterUsed(iLayer, currentCluster.clusterId)) { continue; } - const float meanDeltaR = mTrkParams[iteration].LayerRadii[iLayer + 1] - mTrkParams[iteration].LayerRadii[iLayer]; + const float inverseR0 = 1.f / currentCluster.radius; - for (int iCluster = 0; iCluster < int(layer0.size()); ++iCluster) { - const Cluster& currentCluster = layer0[iCluster]; - const int currentSortedIndex = mTimeFrame->getSortedIndex(targetROF0, iLayer, iCluster); - if (mTimeFrame->isClusterUsed(iLayer, currentCluster.clusterId)) { + for (int iV = startVtx; iV < endVtx; ++iV) { + const auto& pv = primaryVertices[iV]; + if ((pv.isFlagSet(Vertex::Flags::UPCMode) && iteration != 3) || (iteration == 3 && !pv.isFlagSet(Vertex::Flags::UPCMode))) { continue; } - const float inverseR0 = 1.f / currentCluster.radius; - for (int iV = startVtx; iV < endVtx; ++iV) { - const auto& pv = primaryVertices[iV]; - if ((pv.isFlagSet(Vertex::Flags::UPCMode) && iteration != 3) || (iteration == 3 && !pv.isFlagSet(Vertex::Flags::UPCMode))) { - continue; - } + const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(mTimeFrame->getPositionResolution(iLayer)) + math_utils::Sq(mTrkParams[iteration].PVres) / float(pv.getNContributors())); + const float tanLambda = (currentCluster.zCoordinate - pv.getZ()) * inverseR0; + const float zAtRmin = tanLambda * (mTimeFrame->getMinR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate; + const float zAtRmax = tanLambda * (mTimeFrame->getMaxR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate; + const float sqInvDeltaZ0 = 1.f / (math_utils::Sq(currentCluster.zCoordinate - pv.getZ()) + constants::Tolerance); + const float sigmaZ = o2::gpu::CAMath::Sqrt( + math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInvDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * mTimeFrame->getMSangle(iLayer))); - const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(mTimeFrame->getPositionResolution(iLayer)) + math_utils::Sq(mTrkParams[iteration].PVres) / float(pv.getNContributors())); - const float tanLambda = (currentCluster.zCoordinate - pv.getZ()) * inverseR0; - const float zAtRmin = tanLambda * (mTimeFrame->getMinR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate; - const float zAtRmax = tanLambda * (mTimeFrame->getMaxR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate; - const float sqInvDeltaZ0 = 1.f / (math_utils::Sq(currentCluster.zCoordinate - pv.getZ()) + constants::Tolerance); - const float sigmaZ = o2::gpu::CAMath::Sqrt( - math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInvDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * mTimeFrame->getMSangle(iLayer))); + auto bins = getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax, sigmaZ * mTrkParams[iteration].NSigmaCut, mTimeFrame->getPhiCut(iLayer)); + if (bins.x == 0 && bins.y == 0 && bins.z == 0 && bins.w == 0) { + continue; + } + int phiBinsNum = bins.w - bins.y + 1; + if (phiBinsNum < 0) { + phiBinsNum += mTrkParams[iteration].PhiBins; + } - auto bins = getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax, sigmaZ * mTrkParams[iteration].NSigmaCut, mTimeFrame->getPhiCut(iLayer)); - if (bins.x == 0 && bins.y == 0 && bins.z == 0 && bins.w == 0) { + for (int targetROF{minROF}; targetROF <= maxROF; ++targetROF) { + if (!mTimeFrame->mMultiplicityCutMask[targetROF]) { continue; } - int phiBinsNum = bins.w - bins.y + 1; - if (phiBinsNum < 0) { - phiBinsNum += mTrkParams[iteration].PhiBins; + auto layer1 = mTimeFrame->getClustersOnLayer(targetROF, iLayer + 1); + if (layer1.empty()) { + continue; } - - for (int targetROF1{minROF}; targetROF1 <= maxROF; ++targetROF1) { - if (!mTimeFrame->mMultiplicityCutMask[targetROF1] || std::abs(targetROF0 - targetROF1) > mTrkParams[iteration].DeltaROF) { - continue; - } - auto layer1 = mTimeFrame->getClustersOnLayer(targetROF1, iLayer + 1); - if (layer1.empty()) { - continue; - } - for (int iPhi = 0; iPhi < phiBinsNum; ++iPhi) { - int iPhiBin = (bins.y + iPhi) % mTrkParams[iteration].PhiBins; - int firstBinIdx = mTimeFrame->mIndexTableUtils.getBinIndex(bins.x, iPhiBin); - int maxBinIdx = firstBinIdx + (bins.z - bins.x) + 1; - int firstRow = mTimeFrame->getIndexTable(targetROF1, iLayer + 1)[firstBinIdx]; - int lastRow = mTimeFrame->getIndexTable(targetROF1, iLayer + 1)[maxBinIdx]; - for (int iNext = firstRow; iNext < lastRow; ++iNext) { - if (iNext >= int(layer1.size())) { - break; - } - const Cluster& nextCluster = layer1[iNext]; - if (mTimeFrame->isClusterUsed(iLayer + 1, nextCluster.clusterId)) { - continue; - } - float deltaPhi = o2::gpu::GPUCommonMath::Abs(currentCluster.phi - nextCluster.phi); - float deltaZ = o2::gpu::GPUCommonMath::Abs((tanLambda * (nextCluster.radius - currentCluster.radius)) + currentCluster.zCoordinate - nextCluster.zCoordinate); + for (int iPhi = 0; iPhi < phiBinsNum; ++iPhi) { + const int iPhiBin = (bins.y + iPhi) % mTrkParams[iteration].PhiBins; + const int firstBinIdx = mTimeFrame->mIndexTableUtils.getBinIndex(bins.x, iPhiBin); + const int maxBinIdx = firstBinIdx + (bins.z - bins.x) + 1; + const int firstRow = mTimeFrame->getIndexTable(targetROF, iLayer + 1)[firstBinIdx]; + const int lastRow = mTimeFrame->getIndexTable(targetROF, iLayer + 1)[maxBinIdx]; + for (int iNext = firstRow; iNext < lastRow; ++iNext) { + if (iNext >= int(layer1.size())) { + break; + } + const Cluster& nextCluster = layer1[iNext]; + if (mTimeFrame->isClusterUsed(iLayer + 1, nextCluster.clusterId)) { + continue; + } + float deltaPhi = o2::gpu::GPUCommonMath::Abs(currentCluster.phi - nextCluster.phi); + float deltaZ = o2::gpu::GPUCommonMath::Abs((tanLambda * (nextCluster.radius - currentCluster.radius)) + currentCluster.zCoordinate - nextCluster.zCoordinate); #ifdef OPTIMISATION_OUTPUT - MCCompLabel label; - int currentId{currentCluster.clusterId}; - int nextId{nextCluster.clusterId}; - for (auto& lab1 : mTimeFrame->getClusterLabels(iLayer, currentId)) { - for (auto& lab2 : mTimeFrame->getClusterLabels(iLayer + 1, nextId)) { - if (lab1 == lab2 && lab1.isValid()) { - label = lab1; - break; - } - } - if (label.isValid()) { + MCCompLabel label; + int currentId{currentCluster.clusterId}; + int nextId{nextCluster.clusterId}; + for (auto& lab1 : mTimeFrame->getClusterLabels(iLayer, currentId)) { + for (auto& lab2 : mTimeFrame->getClusterLabels(iLayer + 1, nextId)) { + if (lab1 == lab2 && lab1.isValid()) { + label = lab1; break; } } - off << std::format("{}\t{:d}\t{}\t{}\t{}\t{}", iLayer, label.isValid(), (tanLambda * (nextCluster.radius - currentCluster.radius) + currentCluster.zCoordinate - nextCluster.zCoordinate) / sigmaZ, tanLambda, resolution, sigmaZ) << std::endl; + if (label.isValid()) { + break; + } + } + off << std::format("{}\t{:d}\t{}\t{}\t{}\t{}", iLayer, label.isValid(), (tanLambda * (nextCluster.radius - currentCluster.radius) + currentCluster.zCoordinate - nextCluster.zCoordinate) / sigmaZ, tanLambda, resolution, sigmaZ) << std::endl; #endif - if (deltaZ / sigmaZ < mTrkParams[iteration].NSigmaCut && - (deltaPhi < mTimeFrame->getPhiCut(iLayer) || - o2::gpu::GPUCommonMath::Abs(deltaPhi - o2::constants::math::TwoPI) < mTimeFrame->getPhiCut(iLayer))) { - float phi = o2::gpu::GPUCommonMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate); - float tanL = (currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius); - if constexpr (decltype(Tag)::value == PassMode::OnePass::value) { - tracklets.emplace_back(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF1, iLayer + 1, iNext), tanL, phi, targetROF0, targetROF1); - } else if constexpr (decltype(Tag)::value == PassMode::TwoPassCount::value) { - ++localCount; - } else if constexpr (decltype(Tag)::value == PassMode::TwoPassInsert::value) { - const int idx = base + offset++; - tracklets[idx] = Tracklet(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF1, iLayer + 1, iNext), tanL, phi, targetROF0, targetROF1); - } + if (deltaZ / sigmaZ < mTrkParams[iteration].NSigmaCut && + ((deltaPhi < mTimeFrame->getPhiCut(iLayer) || o2::gpu::GPUCommonMath::Abs(deltaPhi - o2::constants::math::TwoPI) < mTimeFrame->getPhiCut(iLayer)))) { + const float phi{o2::gpu::CAMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate)}; + const float tanL = (currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius); + if constexpr (decltype(Tag)::value == PassMode::OnePass::value) { + tracklets.emplace_back(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF, iLayer + 1, iNext), tanL, phi, pivotROF, targetROF); + } else if constexpr (decltype(Tag)::value == PassMode::TwoPassCount::value) { + ++localCount; + } else if constexpr (decltype(Tag)::value == PassMode::TwoPassInsert::value) { + const int idx = base + offset++; + tracklets[idx] = Tracklet(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF, iLayer + 1, iNext), tanL, phi, pivotROF, targetROF); } } } @@ -250,7 +248,10 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROF /// Sort tracklets auto& trkl{mTimeFrame->getTracklets()[iLayer]}; tbb::parallel_sort(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) -> bool { - return a.firstClusterIndex < b.firstClusterIndex || (a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex < b.secondClusterIndex); + if (a.firstClusterIndex != b.firstClusterIndex) { + return a.firstClusterIndex < b.firstClusterIndex; + } + return a.secondClusterIndex < b.secondClusterIndex; }); /// Remove duplicates trkl.erase(std::unique(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) -> bool { @@ -297,7 +298,7 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROF }); } }); -} +} // namespace o2::its template void TrackerTraits::computeLayerCells(const int iteration) @@ -327,7 +328,6 @@ void TrackerTraits::computeLayerCells(const int iteration) for (int iNextTracklet{nextLayerFirstTrackletIndex}; iNextTracklet < nextLayerLastTrackletIndex; ++iNextTracklet) { const Tracklet& nextTracklet{mTimeFrame->getTracklets()[iLayer + 1][iNextTracklet]}; const auto& nextLbl = mTimeFrame->getTrackletsLabel(iLayer + 1)[iNextTracklet]; - bool print = false; if (mTimeFrame->getTracklets()[iLayer + 1][iNextTracklet].firstClusterIndex != nextLayerClusterIndex) { break; } @@ -509,7 +509,8 @@ void TrackerTraits::findCellsNeighbours(const int iteration) const auto& trkl01 = mTimeFrame->getTracklets()[iLayer + 1][currentCellSeed.getSecondTrackletIndex()]; const auto& trkl10 = mTimeFrame->getTracklets()[iLayer + 1][nextCellSeed.getFirstTrackletIndex()]; const auto& trkl11 = mTimeFrame->getTracklets()[iLayer + 2][nextCellSeed.getSecondTrackletIndex()]; - if ((std::max({trkl00.getMaxRof(), trkl01.getMaxRof(), trkl10.getMaxRof(), trkl11.getMaxRof()}) - std::min({trkl00.getMinRof(), trkl01.getMinRof(), trkl10.getMinRof(), trkl10.getMinRof()})) > mTrkParams[0].DeltaROF) { + if ((std::max({trkl00.getMaxRof(), trkl01.getMaxRof(), trkl10.getMaxRof(), trkl11.getMaxRof()}) - + std::min({trkl00.getMinRof(), trkl01.getMinRof(), trkl10.getMinRof(), trkl11.getMinRof()})) > mTrkParams[0].DeltaROF) { continue; } } @@ -657,20 +658,6 @@ void TrackerTraits::processNeighbours(int iLayer, int iLevel, const bou CA_DEBUGGER(failed[0]++); continue; } - if (mTrkParams[0].DeltaROF) { // TODO this has to be improved for the staggering - const auto& trklNeigh = mTimeFrame->getTracklets()[iLayer - 1][neighbourCell.getFirstTrackletIndex()]; - short minRof{std::numeric_limits::max()}, maxRof{std::numeric_limits::min()}; - for (int iLayer{0}; iLayer < mTrkParams[0].NLayers; ++iLayer) { - if (const auto clsId = currentCell.getCluster(iLayer); clsId != constants::UnusedIndex) { - const short clsROF = mTimeFrame->getClusterROF(iLayer, clsId); - minRof = std::min(minRof, clsROF); - maxRof = std::max(maxRof, clsROF); - } - } - if ((std::max(trklNeigh.getMaxRof(), maxRof) - std::min(trklNeigh.getMinRof(), minRof)) > mTrkParams[0].DeltaROF) { - continue; - } - } /// Let's start the fitting procedure CellSeed seed{currentCell}; diff --git a/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx b/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx index be7750964b3e7..201f1d064a632 100644 --- a/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx +++ b/Detectors/ITSMFT/ITS/workflow/src/TrackerSpec.cxx @@ -60,7 +60,7 @@ void TrackerDPL::run(ProcessingContext& pc) mITSTrackingInterface.updateTimeDependentParams(pc); mITSTrackingInterface.run(pc); mTimer.Stop(); - LOGP(info, "CPU Reconstruction time for this TF {} s (cpu), {} s (wall)", mTimer.CpuTime() - cput, mTimer.RealTime() - realt); + LOGP(info, "CPU Reconstruction time for this TF {:.2f} s (cpu), {:.2f} s (wall)", mTimer.CpuTime() - cput, mTimer.RealTime() - realt); } void TrackerDPL::finaliseCCDB(ConcreteDataMatcher& matcher, void* obj)