diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index c87b3d36b9a6a..35393be1e8dd5 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -42,30 +42,30 @@ class TimeFrameGPU final : public TimeFrame void popMemoryStack(const int); void registerHostMemory(const int); void unregisterHostMemory(const int); - void initialise(const int, const TrackingParameters&, const int); - void loadIndexTableUtils(const int); - void loadTrackingFrameInfoDevice(const int, const int); - void createTrackingFrameInfoDeviceArray(const int); - void loadUnsortedClustersDevice(const int, const int); - void createUnsortedClustersDeviceArray(const int, const int = NLayers); - void loadClustersDevice(const int, const int); - void createClustersDeviceArray(const int, const int = NLayers); - void loadClustersIndexTables(const int, const int); - void createClustersIndexTablesArray(const int); - void createUsedClustersDevice(const int, const int); - void createUsedClustersDeviceArray(const int, const int = NLayers); + void initialise(const TrackingParameters&, int maxLayers); + void loadIndexTableUtils(); + void loadTrackingFrameInfoDevice(const int); + void createTrackingFrameInfoDeviceArray(); + void loadUnsortedClustersDevice(const int); + void createUnsortedClustersDeviceArray(const int = NLayers); + void loadClustersDevice(const int); + void createClustersDeviceArray(const int = NLayers); + void loadClustersIndexTables(const int); + void createClustersIndexTablesArray(); + void createUsedClustersDevice(const int); + void createUsedClustersDeviceArray(const int = NLayers); void loadUsedClustersDevice(); - void loadROFrameClustersDevice(const int, const int); - void createROFrameClustersDeviceArray(const int); + void loadROFrameClustersDevice(const int); + void createROFrameClustersDeviceArray(); void loadROFCutMask(const int); - void loadVertices(const int); - void loadROFOverlapTable(const int); - void loadROFVertexLookupTable(const int); - void updateROFVertexLookupTable(const int); + void loadVertices(); + void loadROFOverlapTable(); + void loadROFVertexLookupTable(); + void updateROFVertexLookupTable(); /// - void createTrackletsLUTDevice(const int, const int); - void createTrackletsLUTDeviceArray(const int); + void createTrackletsLUTDevice(bool, const int); + void createTrackletsLUTDeviceArray(); void loadTrackletsDevice(); void loadTrackletsLUTDevice(); void loadCellsDevice(); @@ -74,12 +74,12 @@ class TimeFrameGPU final : public TimeFrame void loadTrackSeedsChi2Device(); void loadTrackSeedsDevice(bounded_vector&); void createTrackletsBuffers(const int); - void createTrackletsBuffersArray(const int); + void createTrackletsBuffersArray(); void createCellsBuffers(const int); - void createCellsBuffersArray(const int); + void createCellsBuffersArray(); void createCellsDevice(); void createCellsLUTDevice(const int); - void createCellsLUTDeviceArray(const int); + void createCellsLUTDeviceArray(); void createNeighboursIndexTablesDevice(const int); void createNeighboursDevice(const unsigned int layer); void createNeighboursLUTDevice(const int, const unsigned int); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index bf004426f9134..fe272f6f8d3bb 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -51,7 +51,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, const int** clustersIndexTables, int** trackletsLUTs, gsl::span trackletsLUTsHost, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, bounded_vector& phiCuts, const float resolutionPV, @@ -82,7 +82,7 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, gsl::span nTracklets, int** trackletsLUTs, gsl::span trackletsLUTsHost, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, bounded_vector& phiCuts, const float resolutionPV, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index b9091eebde377..f7c458fba5479 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -52,10 +52,10 @@ void TimeFrameGPU::allocMem(void** ptr, size_t size, bool extAllocator, } template -void TimeFrameGPU::loadIndexTableUtils(const int iteration) +void TimeFrameGPU::loadIndexTableUtils() { GPUTimer timer("loading indextable utils"); - if (!iteration) { + { GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB); allocMem(reinterpret_cast(&mIndexTableUtilsDevice), sizeof(IndexTableUtilsN), this->hasFrameworkAllocator()); } @@ -64,9 +64,9 @@ void TimeFrameGPU::loadIndexTableUtils(const int iteration) } template -void TimeFrameGPU::createUnsortedClustersDeviceArray(const int iteration, const int maxLayers) +void TimeFrameGPU::createUnsortedClustersDeviceArray(const int maxLayers) { - if (!iteration) { + { GPUTimer timer("creating unsorted clusters array"); allocMem(reinterpret_cast(&mUnsortedClustersDeviceArray), NLayers * sizeof(Cluster*), this->hasFrameworkAllocator()); GPUChkErrS(cudaHostRegister(mUnsortedClustersDevice.data(), NLayers * sizeof(Cluster*), cudaHostRegisterPortable)); @@ -81,9 +81,9 @@ void TimeFrameGPU::createUnsortedClustersDeviceArray(const int iteratio } template -void TimeFrameGPU::loadUnsortedClustersDevice(const int iteration, const int layer) +void TimeFrameGPU::loadUnsortedClustersDevice(const int layer) { - if (!iteration) { + { GPUTimer timer(mGpuStreams[layer], "loading unsorted clusters", layer); GPULog("gpu-transfer: loading {} unsorted clusters on layer {}, for {:.2f} MB.", this->mUnsortedClusters[layer].size(), layer, this->mUnsortedClusters[layer].size() * sizeof(Cluster) / constants::MB); allocMemAsync(reinterpret_cast(&mUnsortedClustersDevice[layer]), this->mUnsortedClusters[layer].size() * sizeof(Cluster), mGpuStreams[layer], this->hasFrameworkAllocator()); @@ -93,9 +93,9 @@ void TimeFrameGPU::loadUnsortedClustersDevice(const int iteration, cons } template -void TimeFrameGPU::createClustersDeviceArray(const int iteration, const int maxLayers) +void TimeFrameGPU::createClustersDeviceArray(const int maxLayers) { - if (!iteration) { + { GPUTimer timer("creating sorted clusters array"); allocMem(reinterpret_cast(&mClustersDeviceArray), NLayers * sizeof(Cluster*), this->hasFrameworkAllocator()); GPUChkErrS(cudaHostRegister(mClustersDevice.data(), NLayers * sizeof(Cluster*), cudaHostRegisterPortable)); @@ -110,9 +110,9 @@ void TimeFrameGPU::createClustersDeviceArray(const int iteration, const } template -void TimeFrameGPU::loadClustersDevice(const int iteration, const int layer) +void TimeFrameGPU::loadClustersDevice(const int layer) { - if (!iteration) { + { GPUTimer timer(mGpuStreams[layer], "loading sorted clusters", layer); GPULog("gpu-transfer: loading {} clusters on layer {}, for {:.2f} MB.", this->mClusters[layer].size(), layer, this->mClusters[layer].size() * sizeof(Cluster) / constants::MB); allocMemAsync(reinterpret_cast(&mClustersDevice[layer]), this->mClusters[layer].size() * sizeof(Cluster), mGpuStreams[layer], this->hasFrameworkAllocator()); @@ -122,9 +122,9 @@ void TimeFrameGPU::loadClustersDevice(const int iteration, const int la } template -void TimeFrameGPU::createClustersIndexTablesArray(const int iteration) +void TimeFrameGPU::createClustersIndexTablesArray() { - if (!iteration) { + { GPUTimer timer("creating clustersindextable array"); allocMem(reinterpret_cast(&mClustersIndexTablesDeviceArray), NLayers * sizeof(int*), this->hasFrameworkAllocator()); GPUChkErrS(cudaHostRegister(mClustersIndexTablesDevice.data(), NLayers * sizeof(int*), cudaHostRegisterPortable)); @@ -139,9 +139,9 @@ void TimeFrameGPU::createClustersIndexTablesArray(const int iteration) } template -void TimeFrameGPU::loadClustersIndexTables(const int iteration, const int layer) +void TimeFrameGPU::loadClustersIndexTables(const int layer) { - if (!iteration) { + { GPUTimer timer(mGpuStreams[layer], "loading sorted clusters", layer); GPULog("gpu-transfer: loading clusters indextable for layer {} with {} elements, for {:.2f} MB.", layer, this->mIndexTables[layer].size(), this->mIndexTables[layer].size() * sizeof(int) / constants::MB); allocMemAsync(reinterpret_cast(&mClustersIndexTablesDevice[layer]), this->mIndexTables[layer].size() * sizeof(int), mGpuStreams[layer], this->hasFrameworkAllocator()); @@ -151,9 +151,9 @@ void TimeFrameGPU::loadClustersIndexTables(const int iteration, const i } template -void TimeFrameGPU::createUsedClustersDeviceArray(const int iteration, const int maxLayers) +void TimeFrameGPU::createUsedClustersDeviceArray(const int maxLayers) { - if (!iteration) { + { GPUTimer timer("creating used clusters flags"); allocMem(reinterpret_cast(&mUsedClustersDeviceArray), NLayers * sizeof(uint8_t*), this->hasFrameworkAllocator()); GPUChkErrS(cudaHostRegister(mUsedClustersDevice.data(), NLayers * sizeof(uint8_t*), cudaHostRegisterPortable)); @@ -168,9 +168,9 @@ void TimeFrameGPU::createUsedClustersDeviceArray(const int iteration, c } template -void TimeFrameGPU::createUsedClustersDevice(const int iteration, const int layer) +void TimeFrameGPU::createUsedClustersDevice(const int layer) { - if (!iteration) { + { GPUTimer timer(mGpuStreams[layer], "creating used clusters flags", layer); GPULog("gpu-transfer: creating {} used clusters flags on layer {}, for {:.2f} MB.", this->mUsedClusters[layer].size(), layer, this->mUsedClusters[layer].size() * sizeof(unsigned char) / constants::MB); allocMemAsync(reinterpret_cast(&mUsedClustersDevice[layer]), this->mUsedClusters[layer].size() * sizeof(unsigned char), mGpuStreams[layer], this->hasFrameworkAllocator()); @@ -190,9 +190,9 @@ void TimeFrameGPU::loadUsedClustersDevice() } template -void TimeFrameGPU::createROFrameClustersDeviceArray(const int iteration) +void TimeFrameGPU::createROFrameClustersDeviceArray() { - if (!iteration) { + { GPUTimer timer("creating ROFrame clusters array"); allocMem(reinterpret_cast(&mROFramesClustersDeviceArray), NLayers * sizeof(int*), this->hasFrameworkAllocator()); GPUChkErrS(cudaHostRegister(mROFramesClustersDevice.data(), NLayers * sizeof(int*), cudaHostRegisterPortable)); @@ -207,9 +207,9 @@ void TimeFrameGPU::createROFrameClustersDeviceArray(const int iteration } template -void TimeFrameGPU::loadROFrameClustersDevice(const int iteration, const int layer) +void TimeFrameGPU::loadROFrameClustersDevice(const int layer) { - if (!iteration) { + { GPUTimer timer(mGpuStreams[layer], "loading ROframe clusters", layer); GPULog("gpu-transfer: loading {} ROframe clusters info on layer {}, for {:.2f} MB.", this->mROFramesClusters[layer].size(), layer, this->mROFramesClusters[layer].size() * sizeof(int) / constants::MB); allocMemAsync(reinterpret_cast(&mROFramesClustersDevice[layer]), this->mROFramesClusters[layer].size() * sizeof(int), mGpuStreams[layer], this->hasFrameworkAllocator()); @@ -219,9 +219,9 @@ void TimeFrameGPU::loadROFrameClustersDevice(const int iteration, const } template -void TimeFrameGPU::createTrackingFrameInfoDeviceArray(const int iteration) +void TimeFrameGPU::createTrackingFrameInfoDeviceArray() { - if (!iteration) { + { GPUTimer timer("creating trackingframeinfo array"); allocMem(reinterpret_cast(&mTrackingFrameInfoDeviceArray), NLayers * sizeof(TrackingFrameInfo*), this->hasFrameworkAllocator()); GPUChkErrS(cudaHostRegister(mTrackingFrameInfoDevice.data(), NLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable)); @@ -236,9 +236,9 @@ void TimeFrameGPU::createTrackingFrameInfoDeviceArray(const int iterati } template -void TimeFrameGPU::loadTrackingFrameInfoDevice(const int iteration, const int layer) +void TimeFrameGPU::loadTrackingFrameInfoDevice(const int layer) { - if (!iteration) { + { GPUTimer timer(mGpuStreams[layer], "loading trackingframeinfo", layer); GPULog("gpu-transfer: loading {} tfinfo on layer {}, for {:.2f} MB.", this->mTrackingFrameInfo[layer].size(), layer, this->mTrackingFrameInfo[layer].size() * sizeof(TrackingFrameInfo) / constants::MB); allocMemAsync(reinterpret_cast(&mTrackingFrameInfoDevice[layer]), this->mTrackingFrameInfo[layer].size() * sizeof(TrackingFrameInfo), mGpuStreams[layer], this->hasFrameworkAllocator()); @@ -250,7 +250,7 @@ void TimeFrameGPU::loadTrackingFrameInfoDevice(const int iteration, con template void TimeFrameGPU::loadROFCutMask(const int iteration) { - if (!iteration || iteration == 3) { // we need to re-load the swapped mult-mask in upc iteration + { GPUTimer timer("loading multiplicity cut mask"); const auto& hostTable = *(this->mROFMask); const auto hostView = hostTable.getView(); @@ -270,9 +270,9 @@ void TimeFrameGPU::loadROFCutMask(const int iteration) } template -void TimeFrameGPU::loadVertices(const int iteration) +void TimeFrameGPU::loadVertices() { - if (!iteration) { + { GPUTimer timer("loading seeding vertices"); GPULog("gpu-transfer: loading {} seeding vertices, for {:.2f} MB.", this->mPrimaryVertices.size(), this->mPrimaryVertices.size() * sizeof(Vertex) / constants::MB); allocMem(reinterpret_cast(&mPrimaryVerticesDevice), this->mPrimaryVertices.size() * sizeof(Vertex), this->hasFrameworkAllocator()); @@ -281,9 +281,9 @@ void TimeFrameGPU::loadVertices(const int iteration) } template -void TimeFrameGPU::loadROFOverlapTable(const int iteration) +void TimeFrameGPU::loadROFOverlapTable() { - if (!iteration) { + { GPUTimer timer("initialising device view of ROFOverlapTable"); const auto& hostTable = this->getROFOverlapTable(); const auto& hostView = this->getROFOverlapTableView(); @@ -305,9 +305,9 @@ void TimeFrameGPU::loadROFOverlapTable(const int iteration) } template -void TimeFrameGPU::loadROFVertexLookupTable(const int iteration) +void TimeFrameGPU::loadROFVertexLookupTable() { - if (!iteration) { + { GPUTimer timer("initialising device view of ROFVertexLookupTable"); const auto& hostTable = this->getROFVertexLookupTable(); const auto& hostView = this->getROFVertexLookupTableView(); @@ -329,10 +329,10 @@ void TimeFrameGPU::loadROFVertexLookupTable(const int iteration) } template -void TimeFrameGPU::updateROFVertexLookupTable(const int iteration) +void TimeFrameGPU::updateROFVertexLookupTable() { const auto& hostTable = this->getROFVertexLookupTable(); - if (!iteration) { + { GPUTimer timer("updating device view of ROFVertexLookupTable"); const auto& hostView = this->getROFVertexLookupTableView(); using TableEntry = ROFVertexLookupTable::TableEntry; @@ -345,19 +345,19 @@ void TimeFrameGPU::updateROFVertexLookupTable(const int iteration) } template -void TimeFrameGPU::createTrackletsLUTDeviceArray(const int iteration) +void TimeFrameGPU::createTrackletsLUTDeviceArray() { - if (!iteration) { + { allocMem(reinterpret_cast(&mTrackletsLUTDeviceArray), (NLayers - 1) * sizeof(int*), this->hasFrameworkAllocator()); } } template -void TimeFrameGPU::createTrackletsLUTDevice(const int iteration, const int layer) +void TimeFrameGPU::createTrackletsLUTDevice(bool allocate, const int layer) { GPUTimer timer(mGpuStreams[layer], "creating tracklets LUTs", layer); const int ncls = this->mClusters[layer].size() + 1; - if (!iteration) { + if (allocate) { GPULog("gpu-allocation: creating tracklets LUT for {} elements on layer {}, for {:.2f} MB.", ncls, layer, ncls * sizeof(int) / constants::MB); allocMemAsync(reinterpret_cast(&mTrackletsLUTDevice[layer]), ncls * sizeof(int), mGpuStreams[layer], this->hasFrameworkAllocator()); GPUChkErrS(cudaMemcpyAsync(&mTrackletsLUTDeviceArray[layer], &mTrackletsLUTDevice[layer], sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[layer].get())); @@ -366,9 +366,9 @@ void TimeFrameGPU::createTrackletsLUTDevice(const int iteration, const } template -void TimeFrameGPU::createTrackletsBuffersArray(const int iteration) +void TimeFrameGPU::createTrackletsBuffersArray() { - if (!iteration) { + { GPUTimer timer("creating tracklet buffers array"); allocMem(reinterpret_cast(&mTrackletsDeviceArray), (NLayers - 1) * sizeof(Tracklet*), this->hasFrameworkAllocator()); } @@ -442,9 +442,9 @@ void TimeFrameGPU::loadCellsDevice() } template -void TimeFrameGPU::createCellsLUTDeviceArray(const int iteration) +void TimeFrameGPU::createCellsLUTDeviceArray() { - if (!iteration) { + { GPUTimer timer("creating cells LUTs array"); allocMem(reinterpret_cast(&mCellsLUTDeviceArray), (NLayers - 2) * sizeof(int*), this->hasFrameworkAllocator()); } @@ -461,9 +461,9 @@ void TimeFrameGPU::createCellsLUTDevice(const int layer) } template -void TimeFrameGPU::createCellsBuffersArray(const int iteration) +void TimeFrameGPU::createCellsBuffersArray() { - if (!iteration) { + { GPUTimer timer("creating cells buffers array"); allocMem(reinterpret_cast(&mCellsDeviceArray), (NLayers - 2) * sizeof(CellSeed*), this->hasFrameworkAllocator()); GPUChkErrS(cudaMemcpy(mCellsDeviceArray, mCellsDevice.data(), mCellsDevice.size() * sizeof(CellSeed*), cudaMemcpyHostToDevice)); @@ -646,12 +646,10 @@ void TimeFrameGPU::popMemoryStack(const int iteration) } template -void TimeFrameGPU::initialise(const int iteration, - const TrackingParameters& trkParam, - const int maxLayers) +void TimeFrameGPU::initialise(const TrackingParameters& trkParam, int maxLayers) { mGpuStreams.resize(NLayers); - o2::its::TimeFrame::initialise(iteration, trkParam, maxLayers, false); + o2::its::TimeFrame::initialise(trkParam, maxLayers); } template diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 0359f2cfb0d03..96482c546e542 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -10,16 +10,11 @@ // or submit itself to any jurisdiction. /// -#include -#include #include -#include "DataFormatsITS/TrackITS.h" - -#include "ITStracking/TrackHelpers.h" #include "ITStrackingGPU/TrackerTraitsGPU.h" #include "ITStrackingGPU/TrackingKernels.h" -#include "ITStracking/Constants.h" +#include "ITStracking/Configuration.h" namespace o2::its { @@ -27,28 +22,33 @@ namespace o2::its template void TrackerTraitsGPU::initialiseTimeFrame(const int iteration) { - mTimeFrameGPU->initialise(iteration, this->mTrkParams[iteration], NLayers); - // on default stream - mTimeFrameGPU->loadVertices(iteration); - // TODO these tables can be put in persistent memory - mTimeFrameGPU->loadROFOverlapTable(iteration); // this can be put in constant memory actually - mTimeFrameGPU->loadROFVertexLookupTable(iteration); - // once the tables are in persistent memory just update the vertex one - // mTimeFrameGPU->updateROFVertexLookupTable(iteration); - mTimeFrameGPU->loadIndexTableUtils(iteration); - mTimeFrameGPU->loadROFCutMask(iteration); - // pinned on host - mTimeFrameGPU->createUsedClustersDeviceArray(iteration); - mTimeFrameGPU->createClustersDeviceArray(iteration); - mTimeFrameGPU->createUnsortedClustersDeviceArray(iteration); - mTimeFrameGPU->createClustersIndexTablesArray(iteration); - mTimeFrameGPU->createTrackingFrameInfoDeviceArray(iteration); - mTimeFrameGPU->createROFrameClustersDeviceArray(iteration); - // device array - mTimeFrameGPU->createTrackletsLUTDeviceArray(iteration); - mTimeFrameGPU->createTrackletsBuffersArray(iteration); - mTimeFrameGPU->createCellsBuffersArray(iteration); - mTimeFrameGPU->createCellsLUTDeviceArray(iteration); + mTimeFrameGPU->initialise(this->mTrkParams[iteration], NLayers); + + if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) { + // on default stream + mTimeFrameGPU->loadVertices(); + // TODO these tables can be put in persistent memory + mTimeFrameGPU->loadROFOverlapTable(); // this can be put in constant memory actually + mTimeFrameGPU->loadROFVertexLookupTable(); + // once the tables are in persistent memory just update the vertex one + // mTimeFrameGPU->updateROFVertexLookupTable(); + mTimeFrameGPU->loadIndexTableUtils(); + // pinned on host + mTimeFrameGPU->createUsedClustersDeviceArray(); + mTimeFrameGPU->createClustersDeviceArray(); + mTimeFrameGPU->createUnsortedClustersDeviceArray(); + mTimeFrameGPU->createClustersIndexTablesArray(); + mTimeFrameGPU->createTrackingFrameInfoDeviceArray(); + mTimeFrameGPU->createROFrameClustersDeviceArray(); + // device array + mTimeFrameGPU->createTrackletsLUTDeviceArray(); + mTimeFrameGPU->createTrackletsBuffersArray(); + mTimeFrameGPU->createCellsBuffersArray(); + mTimeFrameGPU->createCellsLUTDeviceArray(); + } + if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass] || this->mTrkParams[iteration].PassFlags[IterationStep::UseUPCMask]) { + mTimeFrameGPU->loadROFCutMask(iteration); + } // push every create artefact on the stack mTimeFrameGPU->pushMemoryStack(iteration); } @@ -65,22 +65,26 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i { // start by queuing loading needed of two last layers for (int iLayer{NLayers}; iLayer-- > NLayers - 2;) { - mTimeFrameGPU->createUsedClustersDevice(iteration, iLayer); - mTimeFrameGPU->loadClustersDevice(iteration, iLayer); - mTimeFrameGPU->loadClustersIndexTables(iteration, iLayer); - mTimeFrameGPU->loadROFrameClustersDevice(iteration, iLayer); + if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) { + mTimeFrameGPU->createUsedClustersDevice(iLayer); + mTimeFrameGPU->loadClustersDevice(iLayer); + mTimeFrameGPU->loadClustersIndexTables(iLayer); + mTimeFrameGPU->loadROFrameClustersDevice(iLayer); + } mTimeFrameGPU->recordEvent(iLayer); } for (int iLayer{this->mTrkParams[iteration].TrackletsPerRoad()}; iLayer--;) { if (iLayer) { // queue loading data of next layer in parallel, this the copies are overlapping with computation kernels - mTimeFrameGPU->createUsedClustersDevice(iteration, iLayer - 1); - mTimeFrameGPU->loadClustersDevice(iteration, iLayer - 1); - mTimeFrameGPU->loadClustersIndexTables(iteration, iLayer - 1); - mTimeFrameGPU->loadROFrameClustersDevice(iteration, iLayer - 1); + if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) { + mTimeFrameGPU->createUsedClustersDevice(iLayer - 1); + mTimeFrameGPU->loadClustersDevice(iLayer - 1); + mTimeFrameGPU->loadClustersIndexTables(iLayer - 1); + mTimeFrameGPU->loadROFrameClustersDevice(iLayer - 1); + } mTimeFrameGPU->recordEvent(iLayer - 1); } - mTimeFrameGPU->createTrackletsLUTDevice(iteration, iLayer); + mTimeFrameGPU->createTrackletsLUTDevice(this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass], iLayer); mTimeFrameGPU->waitEvent(iLayer, iLayer + 1); // wait stream until all data is available countTrackletsInROFsHandler(mTimeFrameGPU->getDeviceIndexTableUtils(), mTimeFrameGPU->getDeviceROFMaskTableView(), @@ -97,7 +101,7 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i mTimeFrameGPU->getDeviceArrayClustersIndexTables(), mTimeFrameGPU->getDeviceArrayTrackletsLUT(), mTimeFrameGPU->getDeviceTrackletsLUTs(), - iteration, + this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices], this->mTrkParams[iteration].NSigmaCut, mTimeFrameGPU->getPhiCuts(), this->mTrkParams[iteration].PVres, @@ -130,7 +134,7 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i mTimeFrameGPU->getNTracklets(), mTimeFrameGPU->getDeviceArrayTrackletsLUT(), mTimeFrameGPU->getDeviceTrackletsLUTs(), - iteration, + this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices], this->mTrkParams[iteration].NSigmaCut, mTimeFrameGPU->getPhiCuts(), this->mTrkParams[iteration].PVres, @@ -149,15 +153,19 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) { // start by queuing loading needed of three last layers for (int iLayer{NLayers}; iLayer-- > NLayers - 3;) { - mTimeFrameGPU->loadUnsortedClustersDevice(iteration, iLayer); - mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration, iLayer); + if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) { + mTimeFrameGPU->loadUnsortedClustersDevice(iLayer); + mTimeFrameGPU->loadTrackingFrameInfoDevice(iLayer); + } mTimeFrameGPU->recordEvent(iLayer); } for (int iLayer{this->mTrkParams[iteration].CellsPerRoad()}; iLayer--;) { if (iLayer) { - mTimeFrameGPU->loadUnsortedClustersDevice(iteration, iLayer - 1); - mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration, iLayer - 1); + if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) { + mTimeFrameGPU->loadUnsortedClustersDevice(iLayer - 1); + mTimeFrameGPU->loadTrackingFrameInfoDevice(iLayer - 1); + } mTimeFrameGPU->recordEvent(iLayer - 1); } diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 49b8f19d68ea6..f133d73a25867 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -322,7 +322,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel( const int** indexTables, Tracklet** tracklets, int** trackletsLUT, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, const float phiCut, const float resolutionPV, @@ -383,7 +383,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel( if (!vertexLUT.isVertexCompatible(layerIndex, pivotROF, primaryVertex)) { continue; } - if ((primaryVertex.isFlagSet(Vertex::Flags::UPCMode) && iteration != 3) || (iteration == 3 && !primaryVertex.isFlagSet(Vertex::Flags::UPCMode))) { + if (primaryVertex.isFlagSet(Vertex::Flags::UPCMode) != selectUPCVertices) { continue; } @@ -576,7 +576,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, const int** clustersIndexTables, int** trackletsLUTs, gsl::span trackletsLUTsHost, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, bounded_vector& phiCuts, const float resolutionPV, @@ -603,7 +603,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, clustersIndexTables, nullptr, trackletsLUTs, - iteration, + selectUPCVertices, NSigmaCut, phiCuts[layer], resolutionPV, @@ -635,7 +635,7 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, gsl::span nTracklets, int** trackletsLUTs, gsl::span trackletsLUTsHost, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, bounded_vector& phiCuts, const float resolutionPV, @@ -662,7 +662,7 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, clustersIndexTables, tracklets, trackletsLUTs, - iteration, + selectUPCVertices, NSigmaCut, phiCuts[layer], resolutionPV, @@ -1094,7 +1094,7 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const int** clustersIndexTables, int** trackletsLUTs, gsl::span trackletsLUTsHost, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, bounded_vector& phiCuts, const float resolutionPV, @@ -1124,7 +1124,7 @@ template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, gsl::span nTracklets, int** trackletsLUTs, gsl::span trackletsLUTsHost, - const int iteration, + const bool selectUPCVertices, const float NSigmaCut, bounded_vector& phiCuts, const float resolutionPV, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h index dbce5e0dc08a7..ce7b3e5a87630 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h @@ -23,12 +23,25 @@ #include #endif +#include "CommonUtils/EnumFlags.h" #include "DetectorsBase/Propagator.h" #include "ITStracking/Constants.h" namespace o2::its { +// Steering of dedicated steps in an iteration +enum class IterationStep : uint8_t { + FirstPass = 0, + RebuildClusterLUT, + UseUPCMask, + SelectUPCVertices, + ResetVertices, + SkipROFsAboveThreshold, + MarkVerticesAsUPC, +}; +using IterationSteps = o2::utils::EnumFlags; + struct TrackingParameters { int CellMinimumLevel() const noexcept { return MinTrackLength - constants::ClustersPerCell + 1; } int NeighboursPerRoad() const noexcept { return NLayers - 3; } @@ -36,6 +49,7 @@ struct TrackingParameters { int TrackletsPerRoad() const noexcept { return NLayers - 1; } std::string asString() const; + IterationSteps PassFlags{IterationStep::FirstPass, IterationStep::RebuildClusterLUT}; int NLayers = 7; std::vector AddTimeError = {0, 0, 0, 0, 0, 0, 0}; std::vector LayerZ = {16.333f + 1, 16.333f + 1, 16.333f + 1, 42.140f + 1, 42.140f + 1, 73.745f + 1, 73.745f + 1}; @@ -73,9 +87,7 @@ struct TrackingParameters { bool SaveTimeBenchmarks = false; bool DoUPCIteration = false; bool FataliseUponFailure = true; - - bool createArtefactLabels{false}; - + bool CreateArtefactLabels{false}; bool PrintMemory = false; // print allocator usage in epilog report size_t MaxMemory = std::numeric_limits::max(); bool DropTFUponFailure = false; @@ -84,6 +96,7 @@ struct TrackingParameters { struct VertexingParameters { std::string asString() const; + IterationSteps PassFlags{IterationStep::FirstPass, IterationStep::ResetVertices}; 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 diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index 300abb2a3b10d..b78540bddfabf 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -177,7 +177,7 @@ struct TimeFrame { auto& getCellsLabel(int layer) { return mCellLabels[layer]; } bool hasMCinformation() const { return mClusterLabels[0] != nullptr; } - void initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers = NLayers, bool resetVertices = true); + void initialise(const TrackingParameters& trkParam, const int maxLayers = NLayers); bool isClusterUsed(int layer, int clusterId) const { return mUsedClusters[layer][clusterId]; } void markUsedCluster(int layer, int clusterId) { mUsedClusters[layer][clusterId] = true; } diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index 1c3c642429686..aa4592c63f404 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -18,7 +18,6 @@ #include -#include "DetectorsBase/Propagator.h" #include "ITStracking/Configuration.h" #include "ITStracking/IndexTableUtils.h" #include "ITStracking/TimeFrame.h" @@ -46,7 +45,7 @@ class TrackerTraits virtual ~TrackerTraits() = default; virtual void adoptTimeFrame(TimeFrame* tf) { mTimeFrame = tf; } - virtual void initialiseTimeFrame(const int iteration) { mTimeFrame->initialise(iteration, mTrkParams[iteration], mTrkParams[iteration].NLayers, false); } + virtual void initialiseTimeFrame(const int iteration) { mTimeFrame->initialise(mTrkParams[iteration], mTrkParams[iteration].NLayers); } virtual void computeLayerTracklets(const int iteration, int iVertex); virtual void computeLayerCells(const int iteration); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h index 3230737a0f87c..daf8d708e1e23 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h @@ -53,14 +53,11 @@ class VertexerTraits VertexerTraits() = default; virtual ~VertexerTraits() = default; - GPUhd() static const int2 getPhiBins(float phi, float deltaPhi, const IndexTableUtilsN&); - GPUhd() const int2 getPhiBins(float phi, float deltaPhi) { return getPhiBins(phi, deltaPhi, mIndexTableUtils); } - // virtual vertexer interface - virtual void initialise(const TrackingParameters& trackingParams, const int iteration = 0); - virtual void computeTracklets(const int iteration = 0); - virtual void computeTrackletMatching(const int iteration = 0); - virtual void computeVertices(const int iteration = 0); + virtual void initialise(const TrackingParameters& trackingParams); + virtual void computeTracklets(const int iteration); + virtual void computeTrackletMatching(const int iteration); + virtual void computeVertices(const int iteration); virtual void adoptTimeFrame(TimeFrameN* tf) noexcept { mTimeFrame = tf; } virtual void updateVertexingParameters(const std::vector& vrtPar); @@ -115,19 +112,6 @@ class VertexerTraits std::shared_ptr mTaskArena; }; -template -inline void VertexerTraits::initialise(const TrackingParameters& trackingParams, const int iteration) -{ - mTimeFrame->initialise(0, trackingParams, 3, (bool)(!iteration)); // iteration for initialisation must be 0 for correctly resetting the frame, we need to pass the non-reset flag for vertices as well, tho. -} - -template -GPUhdi() const int2 VertexerTraits::getPhiBins(float phi, float dPhi, const IndexTableUtilsN& utils) -{ - return int2{utils.getPhiBinIndex(math_utils::getNormalizedPhi(phi - dPhi)), - utils.getPhiBinIndex(math_utils::getNormalizedPhi(phi + dPhi))}; -} - } // namespace its } // namespace o2 #endif diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx index 49bf9b5b1887d..c425d467a8061 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx @@ -143,7 +143,7 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode // check if something was overridden via configurable params if (ip < constants::MaxIter) { if (tc.startLayerMask[ip] > 0) { - trackParams[2].StartLayerMask = tc.startLayerMask[ip]; + param.StartLayerMask = tc.startLayerMask[ip]; } if (tc.minTrackLgtIter[ip] > 0) { param.MinTrackLength = tc.minTrackLgtIter[ip]; @@ -174,6 +174,14 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode LOGP(fatal, "Unsupported ITS tracking mode {} ", toString(mode)); } + for (auto& param : trackParams) { + param.PassFlags.reset(); + } + trackParams[0].PassFlags.set(IterationStep::FirstPass, IterationStep::RebuildClusterLUT); + if (trackParams.size() > 3 && tc.doUPCIteration) { + trackParams[3].PassFlags.set(IterationStep::UseUPCMask, IterationStep::RebuildClusterLUT, IterationStep::SelectUPCVertices); + } + float bFactor = std::abs(o2::base::Propagator::Instance()->getNominalBz()) / 5.0066791f; float bFactorTracklets = bFactor < 0.01f ? 1.f : bFactor; // for tracklets only @@ -188,7 +196,7 @@ std::vector TrackingMode::getTrackingParameters(TrackingMode p.ReseedIfShorter = tc.reseedIfShorter; p.RepeatRefitOut = tc.repeatRefitOut; p.ShiftRefToCluster = tc.shiftRefToCluster; - p.createArtefactLabels = tc.createArtefactLabels; + p.CreateArtefactLabels = tc.createArtefactLabels; p.PrintMemory = tc.printMemory; p.MaxMemory = tc.maxMemory; @@ -241,6 +249,12 @@ std::vector TrackingMode::getVertexingParameters(TrackingMo { const auto& vc = o2::its::VertexerParamConfig::Instance(); std::vector vertParams(2); // The number of actual iterations will be set as a configKeyVal to allow for pp/PbPb choice + for (auto& param : vertParams) { + param.PassFlags.reset(); + } + vertParams[0].PassFlags.set(IterationStep::FirstPass, IterationStep::ResetVertices); + vertParams[1].PassFlags.set(IterationStep::SkipROFsAboveThreshold, IterationStep::MarkVerticesAsUPC); + // global parameters set for every iteration for (auto& p : vertParams) { p.vertPerRofThreshold = vc.vertPerRofThreshold; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index cafddfcc41a76..fc99bf0f35403 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -241,14 +241,14 @@ void TimeFrame::prepareClusters(const TrackingParameters& trkParam, con } template -void TimeFrame::initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers, bool resetVertices) +void TimeFrame::initialise(const TrackingParameters& trkParam, const int maxLayers) { - if (iteration == 0) { + if (trkParam.PassFlags[IterationStep::FirstPass]) { deepVectorClear(mTracks); deepVectorClear(mTracksLabel); deepVectorClear(mLines); deepVectorClear(mLinesLabels); - if (resetVertices) { + if (trkParam.PassFlags[IterationStep::ResetVertices]) { deepVectorClear(mPrimaryVertices); deepVectorClear(mPrimaryVerticesLabels); } @@ -293,7 +293,7 @@ void TimeFrame::initialise(const int iteration, const TrackingParameter for (auto& v : mNTrackletsPerROF) { v = bounded_vector(getNrof(1) + 1, 0, mMemoryPool.get()); } - if (iteration == 0 || iteration == 3) { + if (trkParam.PassFlags[IterationStep::RebuildClusterLUT]) { prepareClusters(trkParam, maxLayers); } mTotalTracklets = {0, 0}; diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index 382f2314b2e6a..f17d961fc7bb7 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -68,7 +68,7 @@ float Tracker::clustersToTracks(const LogFunc& logger, const LogFunc& e try { for (iteration = 0; iteration < (int)mTrkParams.size(); ++iteration) { mMemoryPool->setMaxMemory(mTrkParams[iteration].MaxMemory); - if (iteration == 3 && mTrkParams[0].DoUPCIteration) { + if (mTrkParams[iteration].PassFlags[IterationStep::UseUPCMask]) { mTimeFrame->useUPCMask(); } float timeFrame{0.}, timeTracklets{0.}, timeCells{0.}, timeNeighbours{0.}, timeRoads{0.}; diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index dc2d6e8889973..19cae4b70f158 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -101,7 +101,7 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer if (!mTimeFrame->getROFVertexLookupTableView().isVertexCompatible(iLayer, pivotROF, pv)) { continue; } - if ((pv.isFlagSet(Vertex::Flags::UPCMode) && iteration != 3) || (iteration == 3 && !pv.isFlagSet(Vertex::Flags::UPCMode))) { + if (pv.isFlagSet(Vertex::Flags::UPCMode) != mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices]) { continue; } const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(mTimeFrame->getPositionResolution(iLayer)) + math_utils::Sq(mTrkParams[iteration].PVres) / float(pv.getNContributors())); @@ -224,7 +224,7 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer }); /// Create tracklets labels - if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].createArtefactLabels) { + if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].CreateArtefactLabels) { tbb::parallel_for(0, mTrkParams[iteration].TrackletsPerRoad(), [&](const int iLayer) { for (auto& trk : mTimeFrame->getTracklets()[iLayer]) { MCCompLabel label; @@ -256,7 +256,7 @@ void TrackerTraits::computeLayerCells(const int iteration) if (iLayer > 0) { deepVectorClear(mTimeFrame->getCellsLookupTable()[iLayer - 1]); } - if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].createArtefactLabels) { + if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].CreateArtefactLabels) { deepVectorClear(mTimeFrame->getCellsLabel(iLayer)); } } @@ -390,7 +390,7 @@ void TrackerTraits::computeLayerCells(const int iteration) std::copy_n(perTrackletCount.begin(), currentLayerTrackletsNum + 1, lut.begin()); } - if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].createArtefactLabels) { + if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].CreateArtefactLabels) { auto& labels = mTimeFrame->getCellsLabel(iLayer); labels.reserve(layerCells.size()); for (const auto& cell : layerCells) { diff --git a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx index 2acbec7fe8bc0..f21dfb40b1ed1 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx @@ -63,7 +63,9 @@ float Vertexer::clustersToVertices(LogFunc logger) logger(fmt::format("=== ITS {} Seeding vertexer iteration {} summary:", mTraits->getName(), iteration)); trkPars.PhiBins = mTraits->getVertexingParameters()[0].PhiBins; trkPars.ZBins = mTraits->getVertexingParameters()[0].ZBins; - auto timeInitIteration = evaluateTask(&Vertexer::initialiseVertexer, StateNames[mCurStep = Init], iteration, evalLog, trkPars, iteration); + trkPars.PassFlags = mVertParams[iteration].PassFlags; + trkPars.PassFlags.set(IterationStep::FirstPass, IterationStep::RebuildClusterLUT); + auto timeInitIteration = evaluateTask(&Vertexer::initialiseVertexer, StateNames[mCurStep = Init], iteration, evalLog, trkPars); auto timeTrackletIteration = evaluateTask(&Vertexer::findTracklets, StateNames[mCurStep = Trackleting], iteration, evalLog, iteration); nTracklets01 = mTimeFrame->getTotalTrackletsTF(0); nTracklets12 = mTimeFrame->getTotalTrackletsTF(1); diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index 00674b715b97d..237e99e57e0da 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -156,6 +156,12 @@ void trackletSelectionKernelHost( } } // namespace +template +void VertexerTraits::initialise(const TrackingParameters& trackingParams) +{ + mTimeFrame->initialise(trackingParams, 3); +} + template void VertexerTraits::updateVertexingParameters(const std::vector& vrtPar) { @@ -528,7 +534,7 @@ void VertexerTraits::computeVertices(const int iteration) cluster.getRMS2(), (ushort)cluster.getSize(), cluster.getAvgDistance2()}; - if (iteration) { + if (mVrtParams[iteration].PassFlags[IterationStep::MarkVerticesAsUPC]) { vertex.setFlags(Vertex::UPCMode); } vertex.setTimeStamp(cluster.getTimeStamp()); @@ -629,7 +635,8 @@ 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; + return mVrtParams[iteration].PassFlags[IterationStep::SkipROFsAboveThreshold] && + (int)mTimeFrame->getROFVertexLookupTableView().getVertices(1, rof).getEntries() > mVrtParams[iteration].vertPerRofThreshold; } template class VertexerTraits<7>; diff --git a/Detectors/Upgrades/ALICE3/TRK/workflow/src/TrackerSpec.cxx b/Detectors/Upgrades/ALICE3/TRK/workflow/src/TrackerSpec.cxx index c9d793a3ec78f..cb4cc3897ae9e 100644 --- a/Detectors/Upgrades/ALICE3/TRK/workflow/src/TrackerSpec.cxx +++ b/Detectors/Upgrades/ALICE3/TRK/workflow/src/TrackerSpec.cxx @@ -177,7 +177,7 @@ std::vector TrackerDPL::createTrackingParamsFromCon // params.UseTrackFollowerMix = paramConfig["UseTrackFollowerMix"].get(); // } if (paramConfig.contains("createArtefactLabels")) { - params.createArtefactLabels = paramConfig["createArtefactLabels"].get(); + params.CreateArtefactLabels = paramConfig["createArtefactLabels"].get(); } if (paramConfig.contains("PrintMemory")) { params.PrintMemory = paramConfig["PrintMemory"].get(); @@ -300,7 +300,7 @@ void TrackerDPL::run(ProcessingContext& pc) const auto trackingLoopStart = std::chrono::steady_clock::now(); for (size_t iter{0}; iter < trackingParams.size(); ++iter) { LOGP(info, "{}", trackingParams[iter].asString()); - timeFrame.initialise(iter, trackingParams[iter], 11, false); + timeFrame.initialise(trackingParams[iter], 11); itsTrackerTraits.computeLayerTracklets(iter, -1); LOGP(info, "Number of tracklets in iteration {}: {}", iter, timeFrame.getNumberOfTracklets()); itsTrackerTraits.computeLayerCells(iter);