diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index a83d9d0d52e8f..6a977f8fef21a 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -60,8 +60,6 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, std::vector& radii, bounded_vector& mulScatAng, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template @@ -93,8 +91,6 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, std::vector& radii, bounded_vector& mulScatAng, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template @@ -113,8 +109,6 @@ void countCellsHandler(const Cluster** sortedClusters, const float cellDeltaTanLambdaSigma, const float nSigmaCut, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template @@ -132,8 +126,6 @@ void computeCellsHandler(const Cluster** sortedClusters, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, const float nSigmaCut, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template @@ -150,8 +142,6 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice, const unsigned int nCellsNext, const int maxCellNeighbours, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Stream& stream); template @@ -167,8 +157,6 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice, const unsigned int nCells, const unsigned int nCellsNext, const int maxCellNeighbours, - const int nBlocks, - const int nThreads, gpu::Stream& stream); int filterCellNeighboursHandler(gpuPair*, @@ -193,9 +181,7 @@ void processNeighboursHandler(const int startLayer, const float maxChi2NDF, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads); + o2::its::ExternalAllocator* alloc); template void countTrackSeedHandler(CellSeed* trackSeeds, @@ -214,9 +200,7 @@ void countTrackSeedHandler(CellSeed* trackSeeds, const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads); + o2::its::ExternalAllocator* alloc); template void computeTrackSeedHandler(CellSeed* trackSeeds, @@ -237,9 +221,7 @@ void computeTrackSeedHandler(CellSeed* trackSeeds, const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads); + o2::its::ExternalAllocator* alloc); } // namespace o2::its #endif // ITSTRACKINGGPU_TRACKINGKERNELS_H_ diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index a9b51580f9be7..bd5e7a8bc59f8 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -627,8 +627,7 @@ constexpr auto makeIterTags(std::index_sequence) { return std::array{makeIterTag()...}; } -// FIXME: we have to be careful that the MaxIter does not diverge from the 4 here! -constexpr auto kIterTags = makeIterTags(std::make_index_sequence<4>{}); +constexpr auto kIterTags = makeIterTags(std::make_index_sequence{}); } // namespace detail template diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 3de2871dd458e..f7a416808fec7 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -18,7 +18,6 @@ #include "ITStrackingGPU/TrackerTraitsGPU.h" #include "ITStrackingGPU/TrackingKernels.h" -#include "ITStracking/TrackingConfigParam.h" #include "ITStracking/Constants.h" namespace o2::its @@ -63,8 +62,6 @@ void TrackerTraitsGPU::adoptTimeFrame(TimeFrame* tf) template void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int iVertex) { - const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); - // start by queuing loading needed of two last layers for (int iLayer{NLayers}; iLayer-- > NLayers - 2;) { mTimeFrameGPU->createUsedClustersDevice(iteration, iLayer); @@ -109,8 +106,6 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i this->mTrkParams[iteration].LayerRadii, mTimeFrameGPU->getMSangles(), mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksLayerTracklets[iteration], - conf.nThreadsLayerTracklets[iteration], mTimeFrameGPU->getStreams()); mTimeFrameGPU->createTrackletsBuffers(iLayer); if (mTimeFrameGPU->getNTracklets()[iLayer] == 0) { @@ -144,8 +139,6 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i this->mTrkParams[iteration].LayerRadii, mTimeFrameGPU->getMSangles(), mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksLayerTracklets[iteration], - conf.nThreadsLayerTracklets[iteration], mTimeFrameGPU->getStreams()); } } @@ -153,8 +146,6 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i template void TrackerTraitsGPU::computeLayerCells(const int iteration) { - auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); - // start by queuing loading needed of three last layers for (int iLayer{NLayers}; iLayer-- > NLayers - 3;) { mTimeFrameGPU->loadUnsortedClustersDevice(iteration, iLayer); @@ -194,8 +185,6 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) this->mTrkParams[iteration].CellDeltaTanLambdaSigma, this->mTrkParams[iteration].NSigmaCut, mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksLayerCells[iteration], - conf.nThreadsLayerCells[iteration], mTimeFrameGPU->getStreams()); mTimeFrameGPU->createCellsBuffers(iLayer); if (mTimeFrameGPU->getNCells()[iLayer] == 0) { @@ -215,8 +204,6 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) this->mTrkParams[iteration].MaxChi2ClusterAttachment, this->mTrkParams[iteration].CellDeltaTanLambdaSigma, this->mTrkParams[iteration].NSigmaCut, - conf.nBlocksLayerCells[iteration], - conf.nThreadsLayerCells[iteration], mTimeFrameGPU->getStreams()); } } @@ -224,8 +211,6 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) template void TrackerTraitsGPU::findCellsNeighbours(const int iteration) { - const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); - for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NeighboursPerRoad(); ++iLayer) { const int currentLayerCellsNum{static_cast(mTimeFrameGPU->getNCells()[iLayer])}; const int nextLayerCellsNum{static_cast(mTimeFrameGPU->getNCells()[iLayer + 1])}; @@ -248,8 +233,6 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) nextLayerCellsNum, 1e2, mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksFindNeighbours[iteration], - conf.nThreadsFindNeighbours[iteration], mTimeFrameGPU->getStream(iLayer)); mTimeFrameGPU->createNeighboursDevice(iLayer); if (mTimeFrameGPU->getNNeighbours()[iLayer] == 0) { @@ -267,8 +250,6 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) currentLayerCellsNum, nextLayerCellsNum, 1e2, - conf.nBlocksFindNeighbours[iteration], - conf.nThreadsFindNeighbours[iteration], mTimeFrameGPU->getStream(iLayer)); mTimeFrameGPU->getArrayNNeighbours()[iLayer] = filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer), mTimeFrameGPU->getDeviceNeighbours(iLayer), @@ -282,7 +263,6 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) template void TrackerTraitsGPU::findRoads(const int iteration) { - auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { const int minimumLayer{startLevel - 1}; bounded_vector> trackSeeds(this->getMemoryPool().get()); @@ -305,9 +285,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[0].MaxChi2NDF, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[0].CorrType, - mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksProcessNeighbours[iteration], - conf.nThreadsProcessNeighbours[iteration]); + mTimeFrameGPU->getFrameworkAllocator()); } // fixme: I don't want to move tracks back and forth, but I need a way to use a thrust::allocator that is aware of our managed memory. if (trackSeeds.empty()) { @@ -334,9 +312,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[0].ShiftRefToCluster, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[0].CorrType, - mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksTracksSeeds[iteration], - conf.nThreadsTracksSeeds[iteration]); + mTimeFrameGPU->getFrameworkAllocator()); mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size()); computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), @@ -356,9 +332,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) this->mTrkParams[0].ShiftRefToCluster, mTimeFrameGPU->getDevicePropagator(), this->mTrkParams[0].CorrType, - mTimeFrameGPU->getFrameworkAllocator(), - conf.nBlocksTracksSeeds[iteration], - conf.nThreadsTracksSeeds[iteration]); + mTimeFrameGPU->getFrameworkAllocator()); mTimeFrameGPU->downloadTrackITSExtDevice(); auto& tracks = mTimeFrameGPU->getTrackITSExt(); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 795b568f6174d..54f92411a3df1 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -813,11 +813,9 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, std::vector& radii, bounded_vector& mulScatAng, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams) { - gpu::computeLayerTrackletsMultiROFKernel<<>>( + gpu::computeLayerTrackletsMultiROFKernel<<<60, 256, 0, streams[layer].get()>>>( utils, rofMask, layer, @@ -874,11 +872,9 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, std::vector& radii, bounded_vector& mulScatAng, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams) { - gpu::computeLayerTrackletsMultiROFKernel<<>>( + gpu::computeLayerTrackletsMultiROFKernel<<<60, 256, 0, streams[layer].get()>>>( utils, rofMask, layer, @@ -909,7 +905,7 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, nTracklets[layer] = unique_end - tracklets_ptr; if (layer) { GPUChkErrS(cudaMemsetAsync(trackletsLUTsHost[layer], 0, (nClusters[layer] + 1) * sizeof(int), streams[layer].get())); - gpu::compileTrackletsLookupTableKernel<<>>( + gpu::compileTrackletsLookupTableKernel<<<60, 256, 0, streams[layer].get()>>>( spanTracklets[layer], trackletsLUTsHost[layer], nTracklets[layer]); @@ -934,11 +930,9 @@ void countCellsHandler( const float cellDeltaTanLambdaSigma, const float nSigmaCut, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams) { - gpu::computeLayerCellsKernel<<>>( + gpu::computeLayerCellsKernel<<<60, 256, 0, streams[layer].get()>>>( sortedClusters, // const Cluster** unsortedClusters, // const Cluster** tfInfo, // const TrackingFrameInfo** @@ -972,11 +966,9 @@ void computeCellsHandler( const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, const float nSigmaCut, - const int nBlocks, - const int nThreads, gpu::Streams& streams) { - gpu::computeLayerCellsKernel<<>>( + gpu::computeLayerCellsKernel<<<60, 256, 0, streams[layer].get()>>>( sortedClusters, // const Cluster** unsortedClusters, // const Cluster** tfInfo, // const TrackingFrameInfo** @@ -1006,11 +998,9 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice, const unsigned int nCellsNext, const int maxCellNeighbours, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Stream& stream) { - gpu::computeLayerCellNeighboursKernel<<>>( + gpu::computeLayerCellNeighboursKernel<<<60, 256, 0, stream.get()>>>( cellsLayersDevice, neighboursLUT, neighboursIndexTable, @@ -1040,11 +1030,9 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice, const unsigned int nCells, const unsigned int nCellsNext, const int maxCellNeighbours, - const int nBlocks, - const int nThreads, gpu::Stream& stream) { - gpu::computeLayerCellNeighboursKernel<<>>( + gpu::computeLayerCellNeighboursKernel<<<60, 256, 0, stream.get()>>>( cellsLayersDevice, neighboursLUT, neighboursIndexTable, @@ -1090,9 +1078,7 @@ void processNeighboursHandler(const int startLayer, const float maxChi2NDF, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads) + o2::its::ExternalAllocator* alloc) { constexpr uint64_t Tag = qStr2Tag("ITS_PNH1"); alloc->pushTagOnStack(Tag); @@ -1101,7 +1087,7 @@ void processNeighboursHandler(const int startLayer, thrust::device_vector> foundSeedsTable(nCells[startLayer] + 1, 0, allocInt); auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator(alloc)).on(gpu::Stream::DefaultStream); - gpu::processNeighboursKernel<<>>( + gpu::processNeighboursKernel<<<60, 256>>>( startLayer, startLevel, allCellSeeds, @@ -1123,7 +1109,7 @@ void processNeighboursHandler(const int startLayer, thrust::device_vector> updatedCellId(foundSeedsTable.back(), 0, allocInt); thrust::device_vector, gpu::TypedAllocator>> updatedCellSeed(foundSeedsTable.back(), allocCellSeed); - gpu::processNeighboursKernel<<>>( + gpu::processNeighboursKernel<<<60, 256>>>( startLayer, startLevel, allCellSeeds, @@ -1155,7 +1141,7 @@ void processNeighboursHandler(const int startLayer, foundSeedsTable.resize(lastCellSeedSize + 1); thrust::fill(nosync_policy, foundSeedsTable.begin(), foundSeedsTable.end(), 0); - gpu::processNeighboursKernel<<>>( + gpu::processNeighboursKernel<<<60, 256>>>( iLayer, --level, allCellSeeds, @@ -1181,7 +1167,7 @@ void processNeighboursHandler(const int startLayer, updatedCellSeed.resize(foundSeeds); thrust::fill(nosync_policy, updatedCellSeed.begin(), updatedCellSeed.end(), CellSeed()); - gpu::processNeighboursKernel<<>>( + gpu::processNeighboursKernel<<<60, 256>>>( iLayer, level, allCellSeeds, @@ -1226,16 +1212,14 @@ void countTrackSeedHandler(CellSeed* trackSeeds, const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads) + o2::its::ExternalAllocator* alloc) { // TODO: the minPts&layerRadii is transfered twice // we should allocate this in constant memory and stop these // small transferes! thrust::device_vector minPts(minPtsHost); thrust::device_vector layerRadii(layerRadiiHost); - gpu::fitTrackSeedsKernel<<>>( + gpu::fitTrackSeedsKernel<<<60, 256>>>( trackSeeds, // CellSeed* foundTrackingFrameInfo, // TrackingFrameInfo** unsortedClusters, // Cluster** @@ -1276,13 +1260,11 @@ void computeTrackSeedHandler(CellSeed* trackSeeds, const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads) + o2::its::ExternalAllocator* alloc) { thrust::device_vector minPts(minPtsHost); thrust::device_vector layerRadii(layerRadiiHost); - gpu::fitTrackSeedsKernel<<>>( + gpu::fitTrackSeedsKernel<<<60, 256>>>( trackSeeds, // CellSeed* foundTrackingFrameInfo, // TrackingFrameInfo** unsortedClusters, // Cluster** @@ -1331,8 +1313,6 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, std::vector& radii, bounded_vector& mulScatAng, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, @@ -1363,8 +1343,6 @@ template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, std::vector& radii, bounded_vector& mulScatAng, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template void countCellsHandler<7>(const Cluster** sortedClusters, @@ -1382,8 +1360,6 @@ template void countCellsHandler<7>(const Cluster** sortedClusters, const float cellDeltaTanLambdaSigma, const float nSigmaCut, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template void computeCellsHandler<7>(const Cluster** sortedClusters, @@ -1400,8 +1376,6 @@ template void computeCellsHandler<7>(const Cluster** sortedClusters, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, const float nSigmaCut, - const int nBlocks, - const int nThreads, gpu::Streams& streams); template void countCellNeighboursHandler<7>(CellSeed<7>** cellsLayersDevice, @@ -1417,8 +1391,6 @@ template void countCellNeighboursHandler<7>(CellSeed<7>** cellsLayersDevice, const unsigned int nCellsNext, const int maxCellNeighbours, o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads, gpu::Stream& stream); template void computeCellNeighboursHandler(CellSeed<7>** cellsLayersDevice, @@ -1433,8 +1405,6 @@ template void computeCellNeighboursHandler(CellSeed<7>** cellsLayersDevice, const unsigned int nCells, const unsigned int nCellsNext, const int maxCellNeighbours, - const int nBlocks, - const int nThreads, gpu::Stream& stream); template void processNeighboursHandler<7>(const int startLayer, @@ -1452,9 +1422,7 @@ template void processNeighboursHandler<7>(const int startLayer, const float maxChi2NDF, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads); + o2::its::ExternalAllocator* alloc); template void countTrackSeedHandler(CellSeed<7>* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, @@ -1472,9 +1440,7 @@ template void countTrackSeedHandler(CellSeed<7>* trackSeeds, const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads); + o2::its::ExternalAllocator* alloc); template void computeTrackSeedHandler(CellSeed<7>* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, @@ -1494,8 +1460,6 @@ template void computeTrackSeedHandler(CellSeed<7>* trackSeeds, const bool shiftRefToCluster, const o2::base::Propagator* propagator, const o2::base::PropagatorF::MatCorrType matCorrType, - o2::its::ExternalAllocator* alloc, - const int nBlocks, - const int nThreads); + o2::its::ExternalAllocator* alloc); } // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index 1f55a95ca0d65..dbce5e0dc08a7 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -84,9 +84,9 @@ struct TrackingParameters { struct VertexingParameters { std::string asString() const; - int nIterations = 1; // Number of vertexing passes to perform std::vector LayerZ = {16.333f + 1, 16.333f + 1, 16.333f + 1, 42.140f + 1, 42.140f + 1, 73.745f + 1, 73.745f + 1}; std::vector LayerRadii = {2.33959f, 3.14076f, 3.91924f, 19.6213f, 24.5597f, 34.388f, 39.3329f}; + int vertPerRofThreshold = 0; // Maximum number of vertices per ROF to trigger second a round int ZBins = 1; int PhiBins = 128; float zCut = -1.f; @@ -100,9 +100,7 @@ struct VertexingParameters { float finalSelectionZCut = -1.f; float duplicateDistance2Cut = -1.f; float tanLambdaCut = -1.f; - float vertNsigmaCut = -1.f; - float vertRadiusSigma = -1.f; - float trackletSigma = -1.f; + float NSigmaCut = -1; float maxZPositionAllowed = -1.f; int clusterContributorsCut = -1; int suppressLowMultDebris = -1; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h index 4b2528b62f057..f8009e3ce8008 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h @@ -31,12 +31,13 @@ constexpr float GB = MB * KB; constexpr bool DoTimeBenchmarks = true; constexpr bool SaveTimeBenchmarks = false; -GPUconstexpr() float Tolerance{1e-12}; // numerical tolerance -GPUconstexpr() int ClustersPerCell{3}; -GPUconstexpr() int UnusedIndex{-1}; -GPUconstexpr() float Resolution{0.0005f}; +GPUconstexpr() float Tolerance = 1e-12; // numerical tolerance +GPUconstexpr() int ClustersPerCell = 3; +GPUconstexpr() int UnusedIndex = -1; +GPUconstexpr() float Resolution = 0.0005f; GPUconstexpr() float Radl = 9.36f; // Radiation length of Si [cm] GPUconstexpr() float Rho = 2.33f; // Density of Si [g/cm^3] +GPUconstexpr() int MaxIter = 4; // Max. supported iterations namespace helpers { diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h index 5ffd55f715a1a..acb55eb1cf993 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h @@ -15,6 +15,7 @@ #include #include "CommonUtils/ConfigurableParam.h" #include "CommonUtils/ConfigurableParamHelper.h" +#include "ITStracking/Constants.h" namespace o2::its { @@ -36,11 +37,9 @@ struct VertexerParamConfig : public o2::conf::ConfigurableParamHelper { - // Use TGeo for mat. budget - static const int MaxIter = 4; static const int MinTrackLength = 4; static const int MaxTrackLength = 7; - bool useMatCorrTGeo = false; // use full geometry to corect for material budget accounting in the fits. Default is to use the material budget LUT. - bool useFastMaterial = false; // use faster material approximation for material budget accounting in the fits. - int addTimeError[7] = {0}; // configure the width of the window in BC to be considered for the tracking. - int minTrackLgtIter[MaxIter] = {}; // minimum track length at each iteration, used only if >0, otherwise use code defaults - uint8_t startLayerMask[MaxIter] = {}; // mask of start layer for this iteration (if >0) - float minPtIterLgt[MaxIter * (MaxTrackLength - MinTrackLength + 1)] = {}; // min.pT for given track length at this iteration, used only if >0, otherwise use code defaults - float sysErrY2[7] = {0}; // systematic error^2 in Y per layer - float sysErrZ2[7] = {0}; // systematic error^2 in Z per layer + + bool useMatCorrTGeo = false; // use full geometry to corect for material budget accounting in the fits. Default is to use the material budget LUT. + bool useFastMaterial = false; // use faster material approximation for material budget accounting in the fits. + int addTimeError[7] = {0}; // configure the width of the window in BC to be considered for the tracking. + int minTrackLgtIter[constants::MaxIter] = {}; // minimum track length at each iteration, used only if >0, otherwise use code defaults + uint8_t startLayerMask[constants::MaxIter] = {}; // mask of start layer for this iteration (if >0) + float minPtIterLgt[constants::MaxIter * (MaxTrackLength - MinTrackLength + 1)] = {}; // min.pT for given track length at this iteration, used only if >0, otherwise use code defaults + float sysErrY2[7] = {0}; // systematic error^2 in Y per layer + float sysErrZ2[7] = {0}; // systematic error^2 in Z per layer float maxChi2ClusterAttachment = -1.f; float maxChi2NDF = -1.f; float nSigmaCut = -1.f; @@ -91,7 +89,7 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper { - static constexpr int MaxIter = TrackerParamConfig::MaxIter; - - /// Set nBlocks/nThreads to summarily override all kernel launch parameters in each iteration. - /// Parameters must start with nBlocks/nThreads. - static constexpr int OverrideValue{-1}; - static constexpr char const* BlocksName = "nBlocks"; - static constexpr char const* ThreadsName = "nThreads"; - int nBlocks = OverrideValue; - int nThreads = OverrideValue; - void maybeOverride() const; - - /// Individual kernel launch parameter for each iteration - int nBlocksLayerTracklets[MaxIter] = {60, 60, 60, 60}; - int nThreadsLayerTracklets[MaxIter] = {256, 256, 256, 256}; - - int nBlocksLayerCells[MaxIter] = {60, 60, 60, 60}; - int nThreadsLayerCells[MaxIter] = {256, 256, 256, 256}; - - int nBlocksFindNeighbours[MaxIter] = {60, 60, 60, 60}; - int nThreadsFindNeighbours[MaxIter] = {256, 256, 256, 256}; - - int nBlocksProcessNeighbours[MaxIter] = {60, 60, 60, 60}; - int nThreadsProcessNeighbours[MaxIter] = {256, 256, 256, 256}; - - int nBlocksTracksSeeds[MaxIter] = {60, 60, 60, 60}; - int nThreadsTracksSeeds[MaxIter] = {256, 256, 256, 256}; - - int nBlocksVtxComputeTracklets[2] = {60, 60}; - int nThreadsVtxComputeTracklets[2] = {256, 256}; - - int nBlocksVtxComputeMatching[2] = {60, 60}; - int nThreadsVtxComputeMatching[2] = {256, 256}; - - O2ParamDef(ITSGpuTrackingParamConfig, "ITSGpuTrackingParam"); -}; - } // namespace o2::its #endif diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h index 77218754dbda3..a045ba1639b13 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h @@ -91,7 +91,7 @@ class Vertexer void printEpilog(LogFunc& logger, const unsigned int trackletN01, const unsigned int trackletN12, - const unsigned selectedN, const unsigned int vertexN, const float initT, + const unsigned selectedN, const unsigned int vertexN, const unsigned int totalVertexN, const float trackletT, const float selecT, const float vertexT); void setNThreads(int n, std::shared_ptr& arena) { mTraits->setNThreads(n, arena); } diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h index 5b609c2fa6c85..1adb09551e326 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h @@ -115,6 +115,8 @@ class VertexerTraits // Frame related quantities TimeFrameN* mTimeFrame = nullptr; // observer ptr private: + bool skipROF(int iteration, int rof) const; + std::shared_ptr mMemoryPool; std::shared_ptr mTaskArena; }; diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx index 6c88b61f2df07..49bf9b5b1887d 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx @@ -136,12 +136,12 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode trackParams[3].TrackletMinPt = 0.1f; trackParams[3].CellDeltaTanLambdaSigma *= 4.; } - for (size_t ip = 0; ip < trackParams.size(); ip++) { + for (int ip = 0; ip < (int)trackParams.size(); ip++) { auto& param = trackParams[ip]; param.ZBins = 64; param.PhiBins = 32; // check if something was overridden via configurable params - if (ip < tc.MaxIter) { + if (ip < constants::MaxIter) { if (tc.startLayerMask[ip] > 0) { trackParams[2].StartLayerMask = tc.startLayerMask[ip]; } @@ -149,7 +149,7 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode param.MinTrackLength = tc.minTrackLgtIter[ip]; } for (int ilg = tc.MaxTrackLength; ilg >= tc.MinTrackLength; ilg--) { - int lslot0 = (tc.MaxTrackLength - ilg), lslot = lslot0 + ip * (tc.MaxTrackLength - tc.MinTrackLength + 1); + int lslot0 = (tc.MaxTrackLength - ilg), lslot = lslot0 + (ip * (tc.MaxTrackLength - tc.MinTrackLength + 1)); if (tc.minPtIterLgt[lslot] > 0.) { param.MinPt[lslot0] = tc.minPtIterLgt[lslot]; } @@ -240,27 +240,15 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode std::vector TrackingMode::getVertexingParameters(TrackingMode::Type mode) { const auto& vc = o2::its::VertexerParamConfig::Instance(); - std::vector vertParams; - if (mode == TrackingMode::Async) { - vertParams.resize(2); // The number of actual iterations will be set as a configKeyVal to allow for pp/PbPb choice - vertParams[1].phiCut = 0.015f; - vertParams[1].tanLambdaCut = 0.015f; - } else if (mode == TrackingMode::Sync) { - vertParams.resize(1); - } else if (mode == TrackingMode::Cosmics) { - vertParams.resize(1); - } else { - LOGP(fatal, "Unsupported ITS vertexing mode {} ", toString(mode)); - } - + std::vector vertParams(2); // The number of actual iterations will be set as a configKeyVal to allow for pp/PbPb choice // global parameters set for every iteration for (auto& p : vertParams) { + p.vertPerRofThreshold = vc.vertPerRofThreshold; p.SaveTimeBenchmarks = vc.saveTimeBenchmarks; p.PrintMemory = vc.printMemory; p.MaxMemory = vc.maxMemory; p.DropTFUponFailure = vc.dropTFUponFailure; - p.nIterations = vc.nIterations; - p.trackletSigma = vc.trackletSigma; + p.NSigmaCut = vc.nSigmaCut; p.maxZPositionAllowed = vc.maxZPositionAllowed; p.clusterContributorsCut = vc.clusterContributorsCut; p.suppressLowMultDebris = vc.suppressLowMultDebris; @@ -270,24 +258,35 @@ std::vector TrackingMode::getVertexingParameters(TrackingMo p.nThreads = vc.nThreads; p.ZBins = vc.ZBins; p.PhiBins = vc.PhiBins; - p.useTruthSeeding = vc.useTruthSeeding; + p.maxTrackletsPerCluster = vc.maxTrackletsPerCluster; + p.zCut = vc.zCut; + p.phiCut = vc.phiCut; + p.pairCut = vc.pairCut; + p.clusterCut = vc.clusterCut; + p.coarseZWindow = vc.coarseZWindow; + p.seedDedupZCut = vc.seedDedupZCut; + p.refitDedupZCut = vc.refitDedupZCut; + p.duplicateZCut = vc.duplicateZCut; + p.finalSelectionZCut = vc.finalSelectionZCut; + p.duplicateDistance2Cut = vc.duplicateDistance2Cut; + p.tanLambdaCut = vc.tanLambdaCut; + } + + if (mode == TrackingMode::Async) { + // relax for UPC iteration + vertParams[1].phiCut = 0.015f; + vertParams[1].tanLambdaCut = 0.015f; + vertParams[1].maxTrackletsPerCluster = 2000; + } else if (mode == TrackingMode::Sync || TrackingMode::Cosmics) { + vertParams.resize(1); + } else { + LOGP(fatal, "Unsupported ITS vertexing mode {} ", toString(mode)); + } + + if (vertParams.size() > vc.nIterations) { + vertParams.resize(vc.nIterations); } - // set for now outside to not disturb status quo - vertParams[0].vertNsigmaCut = vc.vertNsigmaCut; - vertParams[0].vertRadiusSigma = vc.vertRadiusSigma; - vertParams[0].maxTrackletsPerCluster = vc.maxTrackletsPerCluster; - vertParams[0].zCut = vc.zCut; - vertParams[0].phiCut = vc.phiCut; - vertParams[0].pairCut = vc.pairCut; - vertParams[0].clusterCut = vc.clusterCut; - vertParams[0].coarseZWindow = vc.coarseZWindow; - vertParams[0].seedDedupZCut = vc.seedDedupZCut; - vertParams[0].refitDedupZCut = vc.refitDedupZCut; - vertParams[0].duplicateZCut = vc.duplicateZCut; - vertParams[0].finalSelectionZCut = vc.finalSelectionZCut; - vertParams[0].duplicateDistance2Cut = vc.duplicateDistance2Cut; - vertParams[0].tanLambdaCut = vc.tanLambdaCut; return vertParams; } diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index dc032a46213a9..fa881789af296 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -31,12 +31,6 @@ using o2::its::constants::GB; template Tracker::Tracker(TrackerTraits* traits) : mTraits(traits) { - /// Initialise standard configuration with 1 iteration - mTrkParams.resize(1); - if (traits->isGPU()) { - ITSGpuTrackingParamConfig::Instance().maybeOverride(); - ITSGpuTrackingParamConfig::Instance().printKeyValues(true, true); - } } template @@ -46,7 +40,6 @@ void Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& er double total{0}; mTraits->updateTrackingParameters(mTrkParams); - mTimeFrame->updateROFVertexLookupTable(); int maxNvertices{-1}; if (mTrkParams[0].PerPrimaryVertexProcessing) { diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackingConfigParam.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackingConfigParam.cxx index 3101c34d4ab8f..47b5f8ffffdb1 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackingConfigParam.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackingConfigParam.cxx @@ -1,4 +1,4 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// Copyright 2019-2026 CERN and copyright holders of ALICE O2. // See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. // All rights not expressly granted are reserved. // @@ -9,36 +9,6 @@ // granted to it by virtue of its status as an Intergovernmental Organization // or submit itself to any jurisdiction. -#include - -#include "Framework/Logger.h" #include "ITStracking/TrackingConfigParam.h" - O2ParamImpl(o2::its::VertexerParamConfig); O2ParamImpl(o2::its::TrackerParamConfig); -O2ParamImpl(o2::its::ITSGpuTrackingParamConfig); - -namespace o2::its -{ - -void ITSGpuTrackingParamConfig::maybeOverride() const -{ - if (nBlocks == OverrideValue && nThreads == OverrideValue) { - return; - } - const auto name = getName(); - auto members = getDataMembers(); - for (auto member : *members) { - if (!member.name.ends_with(BlocksName) && !member.name.ends_with(ThreadsName)) { - if (nBlocks != OverrideValue && member.name.starts_with(BlocksName) && (member.value != nBlocks)) { - o2::conf::ConfigurableParam::setValue(name, member.name, nBlocks); - } - if (nThreads != OverrideValue && member.name.starts_with(ThreadsName) && (member.value != nThreads)) { - o2::conf::ConfigurableParam::setValue(name, member.name, nThreads); - } - } - } - LOGP(info, "Overwriting gpu threading parameters"); -} // namespace o2::its - -} // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx index eb0841888b03e..fcd9024a74709 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx @@ -209,10 +209,7 @@ void ITSTrackingInterface::run(framework::ProcessingContext& pc) auto clockROFspan = rofsinput[clockLayerId]; auto clockTiming = mTimeFrame->getROFOverlapTableView().getClockLayer(); for (auto iRof{0}; iRof < clockROFspan.size(); ++iRof) { - bounded_vector vtxVecLoc; auto& vtxROF = vertROFvec.emplace_back(clockROFspan[iRof]); - vtxROF.setFirstEntry((int)vertices.size()); - if (mRunVertexer) { auto vtxSpan = mTimeFrame->getPrimaryVertices(clockLayerId, iRof); if (o2::its::TrackerParamConfig::Instance().doUPCIteration) { @@ -231,7 +228,6 @@ void ITSTrackingInterface::run(framework::ProcessingContext& pc) } else { vtxROF.setFlag(o2::itsmft::ROFRecord::VtxStdMode); } - vtxROF.setNEntries((int)vtxSpan.size()); } } diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackingLinkDef.h b/Detectors/ITSMFT/ITS/tracking/src/TrackingLinkDef.h index 0640ff98297b9..46af692fe0c15 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackingLinkDef.h +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackingLinkDef.h @@ -39,9 +39,6 @@ #pragma link C++ class o2::its::TrackerParamConfig + ; #pragma link C++ class o2::conf::ConfigurableParamHelper < o2::its::TrackerParamConfig> + ; -#pragma link C++ class o2::its::ITSGpuTrackingParamConfig + ; -#pragma link C++ class o2::conf::ConfigurableParamHelper < o2::its::ITSGpuTrackingParamConfig> + ; - #pragma link C++ class o2::its::FastMultEstConfig + ; #pragma link C++ class o2::conf::ConfigurableParamHelper < o2::its::FastMultEstConfig> + ; diff --git a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx index 222b4801a5767..cbff174634ec8 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx @@ -60,7 +60,7 @@ float Vertexer::clustersToVertices(LogFunc logger) float timeTracklet{0.f}, timeSelection{0.f}, timeVertexing{0.f}, timeInit{0.f}; try { - for (int iteration = 0; iteration < std::min(mVertParams[0].nIterations, (int)mVertParams.size()); ++iteration) { + for (int iteration = 0; iteration < (int)mVertParams.size(); ++iteration) { mMemoryPool->setMaxMemory(mVertParams[iteration].MaxMemory); unsigned int nTracklets01{0}, nTracklets12{0}; logger(fmt::format("=== ITS {} Seeding vertexer iteration {} summary:", mTraits->getName(), iteration)); @@ -71,12 +71,18 @@ float Vertexer::clustersToVertices(LogFunc logger) nTracklets01 = mTimeFrame->getTotalTrackletsTF(0); nTracklets12 = mTimeFrame->getTotalTrackletsTF(1); auto timeSelectionIteration = evaluateTask(&Vertexer::validateTracklets, StateNames[mCurState = Validating], iteration, evalLog, iteration); + const auto nVerticesBefore = mTimeFrame->getPrimaryVertices().size(); auto timeVertexingIteration = evaluateTask(&Vertexer::findVertices, StateNames[mCurState = Finding], iteration, evalLog, iteration); - printEpilog(logger, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), mTimeFrame->getPrimaryVertices().size(), timeInitIteration, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); + const auto nVerticesAfter = mTimeFrame->getPrimaryVertices().size(); + printEpilog(logger, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), nVerticesAfter - nVerticesBefore, nVerticesAfter, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); timeInit += timeInitIteration; timeTracklet += timeTrackletIteration; timeSelection += timeSelectionIteration; timeVertexing += timeVertexingIteration; + + // update LUT with all currently found vertices so in second iteration we can check vertPerROFThreshold + sortVertices(); + mTimeFrame->updateROFVertexLookupTable(); } } catch (const BoundedMemoryResource::MemoryLimitExceeded& err) { handleException(err); @@ -86,8 +92,6 @@ float Vertexer::clustersToVertices(LogFunc logger) LOGP(fatal, "Uncaught exception!"); } - sortVertices(); - return timeInit + timeTracklet + timeSelection + timeVertexing; } @@ -134,12 +138,12 @@ void Vertexer::adoptTimeFrame(TimeFrameN& tf) template void Vertexer::printEpilog(LogFunc& logger, const unsigned int trackletN01, const unsigned int trackletN12, - const unsigned selectedN, const unsigned int vertexN, const float initT, + const unsigned selectedN, const unsigned int vertexN, const unsigned int totalVertexN, const float trackletT, const float selecT, const float vertexT) { logger(fmt::format(" - {} Vertexer: found {} | {} tracklets in: {} ms", mTraits->getName(), trackletN01, trackletN12, trackletT)); logger(fmt::format(" - {} Vertexer: selected {} tracklets in: {} ms", mTraits->getName(), selectedN, selecT)); - logger(fmt::format(" - {} Vertexer: found {} vertices in: {} ms", mTraits->getName(), vertexN, vertexT)); + logger(fmt::format(" - {} Vertexer: found {} vertices in: {} ms (total: {})", mTraits->getName(), vertexN, vertexT, totalVertexN)); if (mVertParams[0].PrintMemory) { mTimeFrame->printArtefactsMemory(); mMemoryPool->print(); diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index a22d2d6c60990..d0baa65c49147 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -168,15 +168,15 @@ void VertexerTraits::computeTracklets(const int iteration) { mTaskArena->execute([&] { tbb::parallel_for(0, mTimeFrame->getNrof(1), [&](const short pivotRofId) { - bool skipROF = !mTimeFrame->getROFMaskView().isROFEnabled(1, pivotRofId); + bool skip = skipROF(iteration, pivotRofId); const auto& rofRange01 = mTimeFrame->getROFOverlapTableView().getOverlap(1, 0, pivotRofId); for (auto targetRofId = rofRange01.getFirstEntry(); targetRofId < rofRange01.getEntriesBound(); ++targetRofId) { const auto timeErr = mTimeFrame->getROFOverlapTableView().getTimeStamp(0, targetRofId, 1, pivotRofId); trackleterKernelHost( - !skipROF ? mTimeFrame->getClustersOnLayer(targetRofId, 0) : gsl::span(), // Clusters to be matched with the next layer in target rof - !skipROF ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), // Clusters to be matched with the current layer in pivot rof - mTimeFrame->getUsedClustersROF(targetRofId, 0), // Span of the used clusters in the target rof - mTimeFrame->getIndexTable(targetRofId, 0).data(), // Index table to access the data on the next layer in target rof + !skip ? mTimeFrame->getClustersOnLayer(targetRofId, 0) : gsl::span(), // Clusters to be matched with the next layer in target rof + !skip ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), // Clusters to be matched with the current layer in pivot rof + mTimeFrame->getUsedClustersROF(targetRofId, 0), // Span of the used clusters in the target rof + mTimeFrame->getIndexTable(targetRofId, 0).data(), // Index table to access the data on the next layer in target rof mVrtParams[iteration].phiCut, mTimeFrame->getTracklets()[0], // Flat tracklet buffer mTimeFrame->getNTrackletsCluster(pivotRofId, 0), // Span of the number of tracklets per each cluster in pivot rof @@ -191,8 +191,8 @@ void VertexerTraits::computeTracklets(const int iteration) for (auto targetRofId = rofRange12.getFirstEntry(); targetRofId < rofRange12.getEntriesBound(); ++targetRofId) { const auto timeErr = mTimeFrame->getROFOverlapTableView().getTimeStamp(2, targetRofId, 1, pivotRofId); trackleterKernelHost( - !skipROF ? mTimeFrame->getClustersOnLayer(targetRofId, 2) : gsl::span(), - !skipROF ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), + !skip ? mTimeFrame->getClustersOnLayer(targetRofId, 2) : gsl::span(), + !skip ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), mTimeFrame->getUsedClustersROF(targetRofId, 2), mTimeFrame->getIndexTable(targetRofId, 2).data(), mVrtParams[iteration].phiCut, @@ -219,14 +219,14 @@ void VertexerTraits::computeTracklets(const int iteration) } tbb::parallel_for(0, mTimeFrame->getNrof(1), [&](const short pivotRofId) { - bool skipROF = !mTimeFrame->getROFMaskView().isROFEnabled(1, pivotRofId); + bool skip = skipROF(iteration, pivotRofId); const int globalOffsetPivot = mTimeFrame->getSortedStartIndex(pivotRofId, 1); const auto& rofRange01 = mTimeFrame->getROFOverlapTableView().getOverlap(1, 0, pivotRofId); for (auto targetRofId = rofRange01.getFirstEntry(); targetRofId < rofRange01.getEntriesBound(); ++targetRofId) { const auto timeErr = mTimeFrame->getROFOverlapTableView().getTimeStamp(0, targetRofId, 1, pivotRofId); trackleterKernelHost( - !skipROF ? mTimeFrame->getClustersOnLayer(targetRofId, 0) : gsl::span(), - !skipROF ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), + !skip ? mTimeFrame->getClustersOnLayer(targetRofId, 0) : gsl::span(), + !skip ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), mTimeFrame->getUsedClustersROF(targetRofId, 0), mTimeFrame->getIndexTable(targetRofId, 0).data(), mVrtParams[iteration].phiCut, @@ -243,8 +243,8 @@ void VertexerTraits::computeTracklets(const int iteration) for (auto targetRofId = rofRange12.getFirstEntry(); targetRofId < rofRange12.getEntriesBound(); ++targetRofId) { const auto timeErr = mTimeFrame->getROFOverlapTableView().getTimeStamp(2, targetRofId, 1, pivotRofId); trackleterKernelHost( - !skipROF ? mTimeFrame->getClustersOnLayer(targetRofId, 2) : gsl::span(), - !skipROF ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), + !skip ? mTimeFrame->getClustersOnLayer(targetRofId, 2) : gsl::span(), + !skip ? mTimeFrame->getClustersOnLayer(pivotRofId, 1) : gsl::span(), mTimeFrame->getUsedClustersROF(targetRofId, 2), mTimeFrame->getIndexTable(targetRofId, 2).data(), mVrtParams[iteration].phiCut, @@ -293,7 +293,7 @@ void VertexerTraits::computeTrackletMatching(const int iteration) tbb::blocked_range(0, (short)mTimeFrame->getNrof(1)), [&](const tbb::blocked_range& Rofs) { for (short pivotRofId = Rofs.begin(); pivotRofId < Rofs.end(); ++pivotRofId) { - if (mTimeFrame->getFoundTracklets(pivotRofId, 0).empty()) { + if (mTimeFrame->getFoundTracklets(pivotRofId, 0).empty() || skipROF(iteration, pivotRofId)) { continue; } mTimeFrame->getLines(pivotRofId).reserve(mTimeFrame->getNTrackletsCluster(pivotRofId, 0).size()); @@ -330,7 +330,6 @@ void VertexerTraits::computeVertices(const int iteration) const int nRofs = mTimeFrame->getNrof(1); std::vector> rofVertices(nRofs); std::vector> rofLabels(nRofs); - const float nsigmaCut = std::min(mVrtParams[iteration].vertNsigmaCut * mVrtParams[iteration].vertNsigmaCut * (mVrtParams[iteration].vertRadiusSigma * mVrtParams[iteration].vertRadiusSigma + mVrtParams[iteration].trackletSigma * mVrtParams[iteration].trackletSigma), 1.98f); const float pairCut2 = mVrtParams[iteration].pairCut * mVrtParams[iteration].pairCut; const float duplicateZCut = mVrtParams[iteration].duplicateZCut > 0.f ? mVrtParams[iteration].duplicateZCut : std::max(4.f * mVrtParams[iteration].pairCut, 0.5f * mVrtParams[iteration].clusterCut); const float duplicateDistance2Cut = mVrtParams[iteration].duplicateDistance2Cut > 0.f ? mVrtParams[iteration].duplicateDistance2Cut : std::max(16.f * pairCut2, 0.0625f * mVrtParams[iteration].clusterCut * mVrtParams[iteration].clusterCut); @@ -352,6 +351,9 @@ void VertexerTraits::computeVertices(const int iteration) settings.memoryPool = mMemoryPool; const auto processROF = [&](const int rofId) { + if (skipROF(iteration, rofId)) { + return; + } auto& lines = mTimeFrame->getLines(rofId); auto clusters = line_vertexer::buildClusters(std::span{lines.data(), lines.size()}, settings); deepVectorClear(lines); // not needed after @@ -508,7 +510,7 @@ void VertexerTraits::computeVertices(const int iteration) for (const auto sortedId : sortedIndices) { const auto& cluster = clusters[selectedIndices[sortedId]]; const auto beamDistance2 = clusterBeamDistance2(cluster); - if (!(beamDistance2 < nsigmaCut)) { + if (!(beamDistance2 < mVrtParams[iteration].NSigmaCut)) { continue; } if (cluster.getSize() < mVrtParams[iteration].clusterContributorsCut) { @@ -620,5 +622,11 @@ void VertexerTraits::setNThreads(int n, std::shared_ptr +bool VertexerTraits::skipROF(int iteration, int rof) const +{ + return iteration && (int)mTimeFrame->getROFVertexLookupTableView().getVertices(1, rof).getEntries() > mVrtParams[iteration].vertPerRofThreshold; +} + template class VertexerTraits<7>; } // namespace o2::its diff --git a/GPU/Workflow/src/GPUWorkflowITS.cxx b/GPU/Workflow/src/GPUWorkflowITS.cxx index fb27df2ec08b9..ac9834d3eacd1 100644 --- a/GPU/Workflow/src/GPUWorkflowITS.cxx +++ b/GPU/Workflow/src/GPUWorkflowITS.cxx @@ -40,7 +40,6 @@ int32_t GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc if (mNTFs == 1 && pc.services().get().inputTimesliceId == 0) { o2::conf::ConfigurableParam::write(o2::base::NameConf::getConfigOutputFileName(pc.services().get().name, o2::its::VertexerParamConfig::Instance().getName()), o2::its::VertexerParamConfig::Instance().getName()); o2::conf::ConfigurableParam::write(o2::base::NameConf::getConfigOutputFileName(pc.services().get().name, o2::its::TrackerParamConfig::Instance().getName()), o2::its::TrackerParamConfig::Instance().getName()); - o2::conf::ConfigurableParam::write(o2::base::NameConf::getConfigOutputFileName(pc.services().get().name, o2::its::ITSGpuTrackingParamConfig::Instance().getName()), o2::its::ITSGpuTrackingParamConfig::Instance().getName()); } return 0; } diff --git a/prodtests/full-system-test/dpl-workflow.sh b/prodtests/full-system-test/dpl-workflow.sh index e954f6875eb30..52407cc3a4073 100755 --- a/prodtests/full-system-test/dpl-workflow.sh +++ b/prodtests/full-system-test/dpl-workflow.sh @@ -119,7 +119,7 @@ EVE_OPT=" --jsons-folder $EDJSONS_DIR" # ITS vertexing settings if [[ $BEAMTYPE == "pp" || $LIGHTNUCLEI == "1" ]]; then - ITS_CONFIG_KEY+="ITSVertexerParam.phiCut=0.4;ITSVertexerParam.tanLambdaCut=0.17;ITSVertexerParam.pairCut=0.0317563;ITSVertexerParam.clusterCut=0.6640964;ITSVertexerParam.coarseZWindow=0.2049018;ITSVertexerParam.seedDedupZCut=0.0711793;ITSVertexerParam.refitDedupZCut=0.0680009;ITSVertexerParam.duplicateZCut=0.1582193;ITSVertexerParam.finalSelectionZCut=0.1081465;ITSVertexerParam.duplicateDistance2Cut=0.0117033;ITSVertexerParam.clusterContributorsCut=2;ITSVertexerParam.seedMemberRadiusZ=0;ITSVertexerParam.vertNsigmaCut=4.0;ITSVertexerParam.vertRadiusSigma=0.0452309;ITSVertexerParam.trackletSigma=0.0025941;ITSVertexerParam.suppressLowMultDebris=0;" + ITS_CONFIG_KEY+="ITSVertexerParam.phiCut=0.4;ITSVertexerParam.tanLambdaCut=0.17;ITSVertexerParam.pairCut=0.0317563;ITSVertexerParam.clusterCut=0.6640964;ITSVertexerParam.coarseZWindow=0.2049018;ITSVertexerParam.seedDedupZCut=0.0711793;ITSVertexerParam.refitDedupZCut=0.0680009;ITSVertexerParam.duplicateZCut=0.1582193;ITSVertexerParam.finalSelectionZCut=0.1081465;ITSVertexerParam.duplicateDistance2Cut=0.0117033;ITSVertexerParam.clusterContributorsCut=2;ITSVertexerParam.seedMemberRadiusZ=0;ITSVertexerParam.nSigmaCut=0.032841;ITSVertexerParam.suppressLowMultDebris=0;" fi if [[ $CTFINPUT != 1 ]]; then