diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index e2542d841c8bf..00d8dfeba2312 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -34,7 +34,7 @@ class TimeFrameGPU : public TimeFrame using typename TimeFrame::ROFMaskTableN; using typename TimeFrame::TrackingTopologyN; using typename TimeFrame::TrackSeedN; - static constexpr int MaxTransitions = TrackingTopologyN::MaxTransitions; + static constexpr int MaxLinks = TrackingTopologyN::MaxLinks; static constexpr int MaxCells = TrackingTopologyN::MaxCells; static constexpr int MaxStreams = MaxCells > NLayers ? MaxCells : NLayers; @@ -157,7 +157,7 @@ class TimeFrameGPU : public TimeFrame void setDevicePropagator(const o2::base::PropagatorImpl* p) final { this->mPropagatorDevice = p; } // Host-specific getters - gsl::span getNTracklets() { return {mNTracklets.data(), static_cast::size_type>(this->mTrackingTopologyView.nTransitions)}; } + gsl::span getNTracklets() { return {mNTracklets.data(), static_cast::size_type>(this->mTrackingTopologyView.nLinks)}; } gsl::span getNCells() { return {mNCells.data(), static_cast::size_type>(this->mTrackingTopologyView.nCells)}; } auto& getArrayNCells() { return mNCells; } gsl::span getNNeighbours() { return {mNNeighbours.data(), static_cast::size_type>(this->mTrackingTopologyView.nCells)}; } @@ -179,7 +179,7 @@ class TimeFrameGPU : public TimeFrame void allocMem(void**, size_t, bool, int32_t = o2::gpu::GPUMemoryResource::MEMORY_GPU); // Abstract owned and unowned memory allocations on default stream // Host-available device buffer sizes - std::array mNTracklets{}; + std::array mNTracklets{}; std::array mNCells{}; std::array mNNeighbours{}; @@ -205,8 +205,8 @@ class TimeFrameGPU : public TimeFrame const int** mClustersIndexTablesDeviceArray; uint8_t** mUsedClustersDeviceArray; const int** mROFramesClustersDeviceArray; - std::array mTrackletsDevice{}; - std::array mTrackletsLUTDevice{}; + std::array mTrackletsDevice{}; + std::array mTrackletsLUTDevice{}; std::array mCellsLUTDevice{}; std::array mNeighboursLUTDevice{}; @@ -258,7 +258,7 @@ inline std::vector TimeFrameGPU::getClusterSizes() template inline size_t TimeFrameGPU::getNumberOfTracklets() const { - return std::accumulate(mNTracklets.begin(), mNTracklets.begin() + this->mTrackingTopologyView.nTransitions, 0); + return std::accumulate(mNTracklets.begin(), mNTracklets.begin() + this->mTrackingTopologyView.nLinks, 0); } template diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 634ac5217a089..7ad15ec9ddd12 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -40,7 +40,7 @@ class ExternalAllocator; template void countTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const typename ROFOverlapTable::View& rofOverlaps, @@ -58,20 +58,20 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, const bool selectUPCVertices, const float NSigmaCut, const typename TrackingTopology::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minR, std::array& maxR, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams); template void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const typename ROFOverlapTable::View& rofOverlaps, @@ -92,13 +92,13 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const bool selectUPCVertices, const float NSigmaCut, const typename TrackingTopology::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minR, std::array& maxR, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index bc2d90165ee96..614ff53786ebf 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -335,25 +335,25 @@ void TimeFrameGPU::loadTrackingTopologies() GPUTimer timer("initialising device views of TrackingTopology"); const auto& hostTopologies = this->getTrackerTopologies(); mDeviceTrackerTopologyViews.resize(hostTopologies.size()); - using LayerTransition = typename TrackingTopologyN::LayerTransition; + using LayerLink = typename TrackingTopologyN::LayerLink; using CellTopology = typename TrackingTopologyN::CellTopology; using Range = typename TrackingTopologyN::Range; using Id = typename TrackingTopologyN::Id; for (size_t iteration = 0; iteration < hostTopologies.size(); ++iteration) { const auto& topology = hostTopologies[iteration]; - LayerTransition* dTransitions{nullptr}; + LayerLink* dLinks{nullptr}; CellTopology* dCells{nullptr}; - Range* dCellsByFirstTransitionIndex{nullptr}; - Id* dCellsByFirstTransition{nullptr}; - allocMem(reinterpret_cast(&dTransitions), topology.getNTransitions() * sizeof(LayerTransition), this->hasFrameworkAllocator()); + Range* dCellsByFirstLinkIndex{nullptr}; + Id* dCellsByFirstLink{nullptr}; + allocMem(reinterpret_cast(&dLinks), topology.getNLinks() * sizeof(LayerLink), this->hasFrameworkAllocator()); allocMem(reinterpret_cast(&dCells), topology.getNCells() * sizeof(CellTopology), this->hasFrameworkAllocator()); - allocMem(reinterpret_cast(&dCellsByFirstTransitionIndex), topology.getNTransitions() * sizeof(Range), this->hasFrameworkAllocator()); - allocMem(reinterpret_cast(&dCellsByFirstTransition), topology.getNCellsByFirstTransition() * sizeof(Id), this->hasFrameworkAllocator()); - GPUChkErrS(cudaMemcpy(dTransitions, topology.getTransitions().data(), topology.getNTransitions() * sizeof(LayerTransition), cudaMemcpyHostToDevice)); + allocMem(reinterpret_cast(&dCellsByFirstLinkIndex), topology.getNLinks() * sizeof(Range), this->hasFrameworkAllocator()); + allocMem(reinterpret_cast(&dCellsByFirstLink), topology.getNCellsByFirstLink() * sizeof(Id), this->hasFrameworkAllocator()); + GPUChkErrS(cudaMemcpy(dLinks, topology.getLinks().data(), topology.getNLinks() * sizeof(LayerLink), cudaMemcpyHostToDevice)); GPUChkErrS(cudaMemcpy(dCells, topology.getCells().data(), topology.getNCells() * sizeof(CellTopology), cudaMemcpyHostToDevice)); - GPUChkErrS(cudaMemcpy(dCellsByFirstTransitionIndex, topology.getCellsByFirstTransitionIndex().data(), topology.getNTransitions() * sizeof(Range), cudaMemcpyHostToDevice)); - GPUChkErrS(cudaMemcpy(dCellsByFirstTransition, topology.getCellsByFirstTransition().data(), topology.getNCellsByFirstTransition() * sizeof(Id), cudaMemcpyHostToDevice)); - mDeviceTrackerTopologyViews[iteration] = topology.getDeviceView(dTransitions, dCells, dCellsByFirstTransitionIndex, dCellsByFirstTransition); + GPUChkErrS(cudaMemcpy(dCellsByFirstLinkIndex, topology.getCellsByFirstLinkIndex().data(), topology.getNLinks() * sizeof(Range), cudaMemcpyHostToDevice)); + GPUChkErrS(cudaMemcpy(dCellsByFirstLink, topology.getCellsByFirstLink().data(), topology.getNCellsByFirstLink() * sizeof(Id), cudaMemcpyHostToDevice)); + mDeviceTrackerTopologyViews[iteration] = topology.getDeviceView(dLinks, dCells, dCellsByFirstLinkIndex, dCellsByFirstLink); } if (!mDeviceTrackerTopologyViews.empty()) { mDeviceTrackingTopologyView = mDeviceTrackerTopologyViews.front(); @@ -380,7 +380,7 @@ template void TimeFrameGPU::createTrackletsLUTDeviceArray() { { - allocMem(reinterpret_cast(&mTrackletsLUTDeviceArray), MaxTransitions * sizeof(int*), this->hasFrameworkAllocator()); + allocMem(reinterpret_cast(&mTrackletsLUTDeviceArray), MaxLinks * sizeof(int*), this->hasFrameworkAllocator()); } } @@ -388,7 +388,7 @@ template void TimeFrameGPU::createTrackletsLUTDevice(bool allocate, const int layer) { GPUTimer timer(mGpuStreams[layer], "creating tracklets LUTs", layer); - const int fromLayer = this->mTrackingTopologyView.getTransition(layer).fromLayer; + const int fromLayer = this->mTrackingTopologyView.getLink(layer).fromLayer; const int ncls = this->mClusters[fromLayer].size() + 1; if (allocate || mTrackletsLUTDevice[layer] == nullptr) { GPULog("gpu-allocation: creating tracklets LUT for {} elements on layer {}, for {:.2f} MB.", ncls, layer, ncls * sizeof(int) / constants::MB); @@ -403,7 +403,7 @@ void TimeFrameGPU::createTrackletsBuffersArray() { { GPUTimer timer("creating tracklet buffers array"); - allocMem(reinterpret_cast(&mTrackletsDeviceArray), MaxTransitions * sizeof(Tracklet*), this->hasFrameworkAllocator()); + allocMem(reinterpret_cast(&mTrackletsDeviceArray), MaxLinks * sizeof(Tracklet*), this->hasFrameworkAllocator()); } } @@ -412,7 +412,7 @@ void TimeFrameGPU::createTrackletsBuffers(const int layer) { GPUTimer timer(mGpuStreams[layer], "creating tracklet buffers", layer); mNTracklets[layer] = 0; - const int fromLayer = this->mTrackingTopologyView.getTransition(layer).fromLayer; + const int fromLayer = this->mTrackingTopologyView.getLink(layer).fromLayer; GPUChkErrS(cudaMemcpyAsync(&mNTracklets[layer], mTrackletsLUTDevice[layer] + this->mClusters[fromLayer].size(), sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[layer].get())); mGpuStreams[layer].sync(); // ensure number of tracklets is correct GPULog("gpu-transfer: creating tracklets buffer for {} elements on layer {}, for {:.2f} MB.", mNTracklets[layer], layer, mNTracklets[layer] * sizeof(Tracklet) / constants::MB); @@ -491,10 +491,10 @@ template void TimeFrameGPU::createCellsLUTDevice(const int layer) { GPUTimer timer(mGpuStreams[layer], "creating cells LUTs", layer); - const int firstTransition = this->mTrackingTopologyView.getCell(layer).firstTransition; - GPULog("gpu-transfer: creating cell LUT for {} elements on layer {}, for {:.2f} MB.", mNTracklets[firstTransition] + 1, layer, (mNTracklets[firstTransition] + 1) * sizeof(int) / constants::MB); - allocMemAsync(reinterpret_cast(&mCellsLUTDevice[layer]), (mNTracklets[firstTransition] + 1) * sizeof(int), mGpuStreams[layer], this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); - GPUChkErrS(cudaMemsetAsync(mCellsLUTDevice[layer], 0, (mNTracklets[firstTransition] + 1) * sizeof(int), mGpuStreams[layer].get())); + const int firstLink = this->mTrackingTopologyView.getCell(layer).firstLink; + GPULog("gpu-transfer: creating cell LUT for {} elements on layer {}, for {:.2f} MB.", mNTracklets[firstLink] + 1, layer, (mNTracklets[firstLink] + 1) * sizeof(int) / constants::MB); + allocMemAsync(reinterpret_cast(&mCellsLUTDevice[layer]), (mNTracklets[firstLink] + 1) * sizeof(int), mGpuStreams[layer], this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); + GPUChkErrS(cudaMemsetAsync(mCellsLUTDevice[layer], 0, (mNTracklets[firstLink] + 1) * sizeof(int), mGpuStreams[layer].get())); GPUChkErrS(cudaMemcpyAsync(&mCellsLUTDeviceArray[layer], &mCellsLUTDevice[layer], sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[layer].get())); } @@ -515,8 +515,8 @@ void TimeFrameGPU::createCellsBuffers(const int layer) { GPUTimer timer(mGpuStreams[layer], "creating cells buffers"); mNCells[layer] = 0; - const int firstTransition = this->mTrackingTopologyView.getCell(layer).firstTransition; - GPUChkErrS(cudaMemcpyAsync(&mNCells[layer], mCellsLUTDevice[layer] + mNTracklets[firstTransition], sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[layer].get())); + const int firstLink = this->mTrackingTopologyView.getCell(layer).firstLink; + GPUChkErrS(cudaMemcpyAsync(&mNCells[layer], mCellsLUTDevice[layer] + mNTracklets[firstLink], sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[layer].get())); mGpuStreams[layer].sync(); // ensure number of cells is correct GPULog("gpu-transfer: creating cell buffer for {} elements on layer {}, for {:.2f} MB.", mNCells[layer], layer, mNCells[layer] * sizeof(CellSeed) / constants::MB); allocMemAsync(reinterpret_cast(&mCellsDevice[layer]), mNCells[layer] * sizeof(CellSeed), mGpuStreams[layer], this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK)); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index 9db46641f26ae..a19759d0577ec 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -78,16 +78,16 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i mTimeFrameGPU->recordEvent(iLayer); } - for (int transitionId{0}; transitionId < hostTopology.nTransitions; ++transitionId) { - const auto transition = hostTopology.getTransition(transitionId); - mTimeFrameGPU->createTrackletsLUTDevice(this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass], transitionId); - mTimeFrameGPU->waitEvent(transitionId, transition.fromLayer); - mTimeFrameGPU->waitEvent(transitionId, transition.toLayer); + for (int linkId{0}; linkId < hostTopology.nLinks; ++linkId) { + const auto link = hostTopology.getLink(linkId); + mTimeFrameGPU->createTrackletsLUTDevice(this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass], linkId); + mTimeFrameGPU->waitEvent(linkId, link.fromLayer); + mTimeFrameGPU->waitEvent(linkId, link.toLayer); countTrackletsInROFsHandler(mTimeFrameGPU->getDeviceIndexTableUtils(), mTimeFrameGPU->getDeviceROFMaskTableView(), - transitionId, - transition.fromLayer, - transition.toLayer, + linkId, + link.fromLayer, + link.toLayer, mTimeFrameGPU->getDeviceROFOverlapTableView(), mTimeFrameGPU->getDeviceROFVertexLookupTableView(), iVertex, @@ -103,25 +103,25 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices], this->mTrkParams[iteration].NSigmaCut, topology, - mTimeFrameGPU->getTransitionPhiCuts(), + mTimeFrameGPU->getLinkPhiCuts(), this->mTrkParams[iteration].PVres, mTimeFrameGPU->getMinRs(), mTimeFrameGPU->getMaxRs(), mTimeFrameGPU->getPositionResolutions(), this->mTrkParams[iteration].LayerRadii, - mTimeFrameGPU->getTransitionMSAngles(), + mTimeFrameGPU->getLinkMSAngles(), mTimeFrameGPU->getFrameworkAllocator(), mTimeFrameGPU->getStreams()); - mTimeFrameGPU->createTrackletsBuffers(transitionId); - if (mTimeFrameGPU->getNTracklets()[transitionId] == 0) { - mTimeFrameGPU->recordEvent(transitionId); + mTimeFrameGPU->createTrackletsBuffers(linkId); + if (mTimeFrameGPU->getNTracklets()[linkId] == 0) { + mTimeFrameGPU->recordEvent(linkId); continue; } computeTrackletsInROFsHandler(mTimeFrameGPU->getDeviceIndexTableUtils(), mTimeFrameGPU->getDeviceROFMaskTableView(), - transitionId, - transition.fromLayer, - transition.toLayer, + linkId, + link.fromLayer, + link.toLayer, mTimeFrameGPU->getDeviceROFOverlapTableView(), mTimeFrameGPU->getDeviceROFVertexLookupTableView(), iVertex, @@ -140,16 +140,16 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices], this->mTrkParams[iteration].NSigmaCut, topology, - mTimeFrameGPU->getTransitionPhiCuts(), + mTimeFrameGPU->getLinkPhiCuts(), this->mTrkParams[iteration].PVres, mTimeFrameGPU->getMinRs(), mTimeFrameGPU->getMaxRs(), mTimeFrameGPU->getPositionResolutions(), this->mTrkParams[iteration].LayerRadii, - mTimeFrameGPU->getTransitionMSAngles(), + mTimeFrameGPU->getLinkMSAngles(), mTimeFrameGPU->getFrameworkAllocator(), mTimeFrameGPU->getStreams()); - mTimeFrameGPU->recordEvent(transitionId); + mTimeFrameGPU->recordEvent(linkId); } } @@ -168,17 +168,17 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) for (int cellTopologyId{hostTopology.nCells}; cellTopologyId--;) { const auto cellTopology = hostTopology.getCell(cellTopologyId); - const auto first = hostTopology.getTransition(cellTopology.firstTransition); - const auto second = hostTopology.getTransition(cellTopology.secondTransition); - const int currentLayerTrackletsNum{static_cast(mTimeFrameGPU->getNTracklets()[cellTopology.firstTransition])}; - if (!currentLayerTrackletsNum || !mTimeFrameGPU->getNTracklets()[cellTopology.secondTransition]) { + const auto first = hostTopology.getLink(cellTopology.firstLink); + const auto second = hostTopology.getLink(cellTopology.secondLink); + const int currentLayerTrackletsNum{static_cast(mTimeFrameGPU->getNTracklets()[cellTopology.firstLink])}; + if (!currentLayerTrackletsNum || !mTimeFrameGPU->getNTracklets()[cellTopology.secondLink]) { mTimeFrameGPU->getNCells()[cellTopologyId] = 0; continue; } mTimeFrameGPU->createCellsLUTDevice(cellTopologyId); - mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.firstTransition); - mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.secondTransition); + mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.firstLink); + mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.secondLink); mTimeFrameGPU->waitEvent(cellTopologyId, first.fromLayer); mTimeFrameGPU->waitEvent(cellTopologyId, first.toLayer); mTimeFrameGPU->waitEvent(cellTopologyId, second.toLayer); @@ -249,7 +249,7 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) for (int sourceCellTopologyId{0}; sourceCellTopologyId < hostTopology.nCells; ++sourceCellTopologyId) { const auto sourceCellTopology = hostTopology.getCell(sourceCellTopologyId); const int sourceCellsNum{static_cast(mTimeFrameGPU->getNCells()[sourceCellTopologyId])}; - if (!sourceCellsNum || sourceCellTopology.secondTransition != targetCellTopology.firstTransition) { + if (!sourceCellsNum || sourceCellTopology.secondLink != targetCellTopology.firstLink) { continue; } mTimeFrameGPU->waitEvent(targetCellTopologyId, sourceCellTopologyId); @@ -279,7 +279,7 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) for (int sourceCellTopologyId{0}; sourceCellTopologyId < hostTopology.nCells; ++sourceCellTopologyId) { const auto sourceCellTopology = hostTopology.getCell(sourceCellTopologyId); const int sourceCellsNum{static_cast(mTimeFrameGPU->getNCells()[sourceCellTopologyId])}; - if (!sourceCellsNum || sourceCellTopology.secondTransition != targetCellTopology.firstTransition) { + if (!sourceCellsNum || sourceCellTopology.secondLink != targetCellTopology.firstLink) { continue; } computeCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(), diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 57f2621b6d9ec..cdeb95c67be30 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -313,8 +313,8 @@ GPUg() void __launch_bounds__(constants::GPUThreads, 1) computeLayerCellsKernel( const float nSigmaCut) { const auto cellTopology = topology.getCell(cellTopologyId); - const auto first = topology.getTransition(cellTopology.firstTransition); - const auto second = topology.getTransition(cellTopology.secondTransition); + const auto first = topology.getLink(cellTopology.firstLink); + const auto second = topology.getLink(cellTopology.secondLink); const int layers[3] = {first.fromLayer, first.toLayer, second.toLayer}; for (int iCurrentTrackletIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackletIndex < nTrackletsCurrent; iCurrentTrackletIndex += blockDim.x * gridDim.x) { if constexpr (!initRun) { @@ -322,19 +322,19 @@ GPUg() void __launch_bounds__(constants::GPUThreads, 1) computeLayerCellsKernel( continue; } } - const Tracklet& currentTracklet = tracklets[cellTopology.firstTransition][iCurrentTrackletIndex]; + const Tracklet& currentTracklet = tracklets[cellTopology.firstLink][iCurrentTrackletIndex]; const int nextLayerClusterIndex{currentTracklet.secondClusterIndex}; - const int nextLayerFirstTrackletIndex{trackletsLUT[cellTopology.secondTransition][nextLayerClusterIndex]}; - const int nextLayerLastTrackletIndex{trackletsLUT[cellTopology.secondTransition][nextLayerClusterIndex + 1]}; + const int nextLayerFirstTrackletIndex{trackletsLUT[cellTopology.secondLink][nextLayerClusterIndex]}; + const int nextLayerLastTrackletIndex{trackletsLUT[cellTopology.secondLink][nextLayerClusterIndex + 1]}; if (nextLayerFirstTrackletIndex == nextLayerLastTrackletIndex) { continue; } int foundCells{0}; for (int iNextTrackletIndex{nextLayerFirstTrackletIndex}; iNextTrackletIndex < nextLayerLastTrackletIndex; ++iNextTrackletIndex) { - if (tracklets[cellTopology.secondTransition][iNextTrackletIndex].firstClusterIndex != nextLayerClusterIndex) { + if (tracklets[cellTopology.secondLink][iNextTrackletIndex].firstClusterIndex != nextLayerClusterIndex) { break; } - const Tracklet& nextTracklet = tracklets[cellTopology.secondTransition][iNextTrackletIndex]; + const Tracklet& nextTracklet = tracklets[cellTopology.secondLink][iNextTrackletIndex]; if (!currentTracklet.getTimeStamp().isCompatible(nextTracklet.getTimeStamp())) { continue; } @@ -396,7 +396,7 @@ template GPUg() void __launch_bounds__(constants::GPUThreads, 1) computeLayerTrackletsMultiROFKernel( const IndexTableUtils* utils, const typename ROFMaskTable::View rofMask, - const int transitionId, + const int linkId, const typename TrackingTopology::View topology, const typename ROFOverlapTable::View rofOverlaps, const typename ROFVertexLookupTable::View vertexLUT, @@ -419,9 +419,9 @@ GPUg() void __launch_bounds__(constants::GPUThreads, 1) computeLayerTrackletsMul const float meanDeltaR, const float MSAngle) { - const auto transition = topology.getTransition(transitionId); - const int fromLayer = transition.fromLayer; - const int toLayer = transition.toLayer; + const auto link = topology.getLink(linkId); + const int fromLayer = link.fromLayer; + const int toLayer = link.toLayer; const int phiBins{utils->getNphiBins()}; const int zBins{utils->getNzBins()}; const int tableSize{phiBins * zBins + 1}; @@ -462,7 +462,7 @@ GPUg() void __launch_bounds__(constants::GPUThreads, 1) computeLayerTrackletsMul continue; } if constexpr (!initRun) { - if (trackletsLUT[transitionId][currentSortedIndex] == trackletsLUT[transitionId][currentSortedIndex + 1]) { + if (trackletsLUT[linkId][currentSortedIndex] == trackletsLUT[linkId][currentSortedIndex + 1]) { continue; } } @@ -523,12 +523,12 @@ GPUg() void __launch_bounds__(constants::GPUThreads, 1) computeLayerTrackletsMul const float deltaZ{o2::gpu::CAMath::Abs(tanLambda * (nextCluster.radius - currentCluster.radius) + currentCluster.zCoordinate - nextCluster.zCoordinate)}; if (deltaZ / sigmaZ < NSigmaCut && (deltaPhi < phiCut || o2::gpu::CAMath::Abs(deltaPhi - o2::constants::math::TwoPI) < phiCut)) { if constexpr (initRun) { - trackletsLUT[transitionId][currentSortedIndex]++; // we need l0 as well for usual exclusive sums. + trackletsLUT[linkId][currentSortedIndex]++; // we need l0 as well for usual exclusive sums. } else { const float phi{o2::gpu::CAMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate)}; const float tanL{(currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius)}; const int nextSortedIndex{ROFClusters[toLayer][targetROF] + nextClusterIndex}; - new (tracklets[transitionId] + trackletsLUT[transitionId][currentSortedIndex] + storedTracklets) Tracklet{currentSortedIndex, nextSortedIndex, tanL, phi, ts}; + new (tracklets[linkId] + trackletsLUT[linkId][currentSortedIndex] + storedTracklets) Tracklet{currentSortedIndex, nextSortedIndex, tanL, phi, ts}; } ++storedTracklets; } @@ -672,7 +672,7 @@ GPUg() void __launch_bounds__(constants::GPUThreads, 1) processNeighboursKernel( template void countTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const typename ROFOverlapTable::View& rofOverlaps, @@ -690,20 +690,20 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, const bool selectUPCVertices, const float NSigmaCut, const typename TrackingTopology::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minRs, std::array& maxRs, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams) { - gpu::computeLayerTrackletsMultiROFKernel<<>>( + gpu::computeLayerTrackletsMultiROFKernel<<>>( utils, rofMask, - transitionId, + linkId, topology, rofOverlaps, vertexLUT, @@ -718,21 +718,21 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, trackletsLUTs, selectUPCVertices, NSigmaCut, - transitionPhiCuts[transitionId], + linkPhiCuts[linkId], resolutionPV, minRs[toLayer], maxRs[toLayer], resolutions[fromLayer], radii[toLayer] - radii[fromLayer], - transitionMSAngles[transitionId]); - auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator(alloc)).on(streams[transitionId].get()); - thrust::exclusive_scan(nosync_policy, trackletsLUTsHost[transitionId], trackletsLUTsHost[transitionId] + nClusters[fromLayer] + 1, trackletsLUTsHost[transitionId]); + linkMSAngles[linkId]); + auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator(alloc)).on(streams[linkId].get()); + thrust::exclusive_scan(nosync_policy, trackletsLUTsHost[linkId], trackletsLUTsHost[linkId] + nClusters[fromLayer] + 1, trackletsLUTsHost[linkId]); } template void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const typename ROFMaskTable::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const typename ROFOverlapTable::View& rofOverlaps, @@ -753,20 +753,20 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const bool selectUPCVertices, const float NSigmaCut, const typename TrackingTopology::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minRs, std::array& maxRs, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams) { - gpu::computeLayerTrackletsMultiROFKernel<<>>( + gpu::computeLayerTrackletsMultiROFKernel<<>>( utils, rofMask, - transitionId, + linkId, topology, rofOverlaps, vertexLUT, @@ -781,25 +781,25 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, trackletsLUTs, selectUPCVertices, NSigmaCut, - transitionPhiCuts[transitionId], + linkPhiCuts[linkId], resolutionPV, minRs[toLayer], maxRs[toLayer], resolutions[fromLayer], radii[toLayer] - radii[fromLayer], - transitionMSAngles[transitionId]); - thrust::device_ptr tracklets_ptr(spanTracklets[transitionId]); - auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator(alloc)).on(streams[transitionId].get()); - thrust::sort(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[transitionId]); - auto unique_end = thrust::unique(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[transitionId]); - nTracklets[transitionId] = unique_end - tracklets_ptr; + linkMSAngles[linkId]); + thrust::device_ptr tracklets_ptr(spanTracklets[linkId]); + auto nosync_policy = THRUST_NAMESPACE::par_nosync(gpu::TypedAllocator(alloc)).on(streams[linkId].get()); + thrust::sort(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[linkId]); + auto unique_end = thrust::unique(nosync_policy, tracklets_ptr, tracklets_ptr + nTracklets[linkId]); + nTracklets[linkId] = unique_end - tracklets_ptr; if (fromLayer > 0) { - GPUChkErrS(cudaMemsetAsync(trackletsLUTsHost[transitionId], 0, (nClusters[fromLayer] + 1) * sizeof(int), streams[transitionId].get())); - gpu::compileTrackletsLookupTableKernel<<>>( - spanTracklets[transitionId], - trackletsLUTsHost[transitionId], - nTracklets[transitionId]); - thrust::exclusive_scan(nosync_policy, trackletsLUTsHost[transitionId], trackletsLUTsHost[transitionId] + nClusters[fromLayer] + 1, trackletsLUTsHost[transitionId]); + GPUChkErrS(cudaMemsetAsync(trackletsLUTsHost[linkId], 0, (nClusters[fromLayer] + 1) * sizeof(int), streams[linkId].get())); + gpu::compileTrackletsLookupTableKernel<<>>( + spanTracklets[linkId], + trackletsLUTsHost[linkId], + nTracklets[linkId]); + thrust::exclusive_scan(nosync_policy, trackletsLUTsHost[linkId], trackletsLUTsHost[linkId] + nClusters[fromLayer] + 1, trackletsLUTsHost[linkId]); } } @@ -1245,7 +1245,7 @@ void computeTrackSeedHandler(TrackSeed* trackSeeds, /// Explicit instantiation of ITS2 handlers template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const ROFMaskTable<7>::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const ROFOverlapTable<7>::View& rofOverlaps, @@ -1263,19 +1263,19 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const bool selectUPCVertices, const float NSigmaCut, const TrackingTopology<7>::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minRs, std::array& maxRs, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams); template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const ROFMaskTable<7>::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const ROFOverlapTable<7>::View& rofOverlaps, @@ -1296,13 +1296,13 @@ template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const bool selectUPCVertices, const float NSigmaCut, const TrackingTopology<7>::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minRs, std::array& maxRs, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams); @@ -1445,7 +1445,7 @@ template void computeTrackSeedHandler(TrackSeed<7>* trackSeeds, #ifdef ENABLE_UPGRADES template void countTrackletsInROFsHandler<11>(const IndexTableUtils<11>* utils, const ROFMaskTable<11>::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const ROFOverlapTable<11>::View& rofOverlaps, @@ -1463,19 +1463,19 @@ template void countTrackletsInROFsHandler<11>(const IndexTableUtils<11>* utils, const bool selectUPCVertices, const float NSigmaCut, const TrackingTopology<11>::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minRs, std::array& maxRs, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams); template void computeTrackletsInROFsHandler<11>(const IndexTableUtils<11>* utils, const ROFMaskTable<11>::View& rofMask, - const int transitionId, + const int linkId, const int fromLayer, const int toLayer, const ROFOverlapTable<11>::View& rofOverlaps, @@ -1496,13 +1496,13 @@ template void computeTrackletsInROFsHandler<11>(const IndexTableUtils<11>* utils const bool selectUPCVertices, const float NSigmaCut, const TrackingTopology<11>::View topology, - bounded_vector& transitionPhiCuts, + bounded_vector& linkPhiCuts, const float resolutionPV, std::array& minRs, std::array& maxRs, bounded_vector& resolutions, std::vector& radii, - bounded_vector& transitionMSAngles, + bounded_vector& linkMSAngles, o2::its::ExternalAllocator* alloc, gpu::Streams& streams); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index 9a31e4014bff5..ae466a32bfc89 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -114,10 +114,10 @@ struct TimeFrame { auto& getMaxRs() { return mMaxR; } float getMinR(int layer) const { return mMinR[layer]; } float getMaxR(int layer) const { return mMaxR[layer]; } - float getTransitionPhiCut(int transitionId) const { return mTransitionPhiCuts[transitionId]; } - float getTransitionMSAngle(int transitionId) const { return mTransitionMSAngles[transitionId]; } - auto& getTransitionPhiCuts() { return mTransitionPhiCuts; } - auto& getTransitionMSAngles() { return mTransitionMSAngles; } + float getLinkPhiCut(int linkId) const { return mLinkPhiCuts[linkId]; } + float getLinkMSAngle(int linkId) const { return mLinkMSAngles[linkId]; } + auto& getLinkPhiCuts() { return mLinkPhiCuts; } + auto& getLinkMSAngles() { return mLinkMSAngles; } float getPositionResolution(int layer) const { return mPositionResolution[layer]; } auto& getPositionResolutions() { return mPositionResolution; } @@ -307,8 +307,8 @@ struct TimeFrame { bool isBeamPositionOverridden = false; std::array mMinR; std::array mMaxR; - bounded_vector mTransitionPhiCuts; - bounded_vector mTransitionMSAngles; + bounded_vector mLinkPhiCuts; + bounded_vector mLinkMSAngles; bounded_vector mPositionResolution; std::array, NLayers> mClusterSize; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h index 2afb67609664f..a9b240d2ef037 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingTopology.h @@ -38,26 +38,26 @@ class TrackingTopology using Id = uint8_t; using Mask = LayerMask; using Range = o2::dataformats::RangeReference; - static constexpr int MaxTransitions = (NLayers * (NLayers - 1)) / 2; + static constexpr int MaxLinks = (NLayers * (NLayers - 1)) / 2; static constexpr int MaxCells = (NLayers * (NLayers - 1) * (NLayers - 2)) / 6; static_assert(NLayers < std::numeric_limits::max()); - static_assert(MaxTransitions <= std::numeric_limits::max()); + static_assert(MaxLinks <= std::numeric_limits::max()); static_assert(MaxCells <= std::numeric_limits::max()); // Describes from which layer to which layer the look-up happens - struct LayerTransition { + struct LayerLink { Id fromLayer{0}; Id toLayer{0}; }; - static_assert(std::is_standard_layout_v); - static_assert(std::is_trivially_copyable_v); - static_assert(sizeof(LayerTransition) == (2 * sizeof(Id))); + static_assert(std::is_standard_layout_v); + static_assert(std::is_trivially_copyable_v); + static_assert(sizeof(LayerLink) == (2 * sizeof(Id))); - // Describes from which LayerTransition a tracklet is allowed to originate - // and with which LayerTransition this can be combined additionally the hitMasked is cached + // Describes from which LayerLink a tracklet is allowed to originate + // and with which LayerLink this can be combined additionally the hitMasked is cached struct CellTopology { - Id firstTransition{0}; - Id secondTransition{0}; + Id firstLink{0}; + Id secondLink{0}; Mask hitLayerMask{0}; }; static_assert(std::is_standard_layout_v); @@ -66,33 +66,33 @@ class TrackingTopology // GPU ready view of the underlying LUTs struct View { - const LayerTransition* transitions{nullptr}; + const LayerLink* links{nullptr}; const CellTopology* cells{nullptr}; - const Range* cellsByFirstTransitionIndex{nullptr}; - const Id* cellsByFirstTransition{nullptr}; - Id nTransitions{0}; + const Range* cellsByFirstLinkIndex{nullptr}; + const Id* cellsByFirstLink{nullptr}; + Id nLinks{0}; Id nCells{0}; - Id nCellsByFirstTransition{0}; + Id nCellsByFirstLink{0}; - GPUhdi() const LayerTransition& getTransition(Id id) const { return transitions[id]; } + GPUhdi() const LayerLink& getLink(Id id) const { return links[id]; } GPUhdi() const CellTopology& getCell(Id id) const { return cells[id]; } - GPUhdi() Range getCellsStartingWithTransition(Id transitionId) const { return cellsByFirstTransitionIndex[transitionId]; } + GPUhdi() Range getCellsStartingWithLink(Id linkId) const { return cellsByFirstLinkIndex[linkId]; } #ifndef GPUCA_GPUCODE std::string asString() const { - std::string out = fmt::format("TrackingTopology: transitions={} cells={}", nTransitions, nCells); - out += "\n transitions:"; - for (Id transitionId = 0; transitionId < nTransitions; ++transitionId) { - const auto& t = transitions[transitionId]; - out += fmt::format("\n {}: {} -> {}", transitionId, t.fromLayer, t.toLayer); + std::string out = fmt::format("TrackingTopology: links={} cells={}", nLinks, nCells); + out += "\n links:"; + for (Id linkId = 0; linkId < nLinks; ++linkId) { + const auto& t = links[linkId]; + out += fmt::format("\n {}: {} -> {}", linkId, t.fromLayer, t.toLayer); } out += "\n cells:"; for (Id cellId = 0; cellId < nCells; ++cellId) { const auto& c = cells[cellId]; - const auto& first = transitions[c.firstTransition]; - const auto& second = transitions[c.secondTransition]; - out += fmt::format("\n {}: {} -> {} -> {} hitMask={} transitions=({}, {})", cellId, first.fromLayer, first.toLayer, second.toLayer, c.hitLayerMask.asString(), c.firstTransition, c.secondTransition); + const auto& first = links[c.firstLink]; + const auto& second = links[c.secondLink]; + out += fmt::format("\n {}: {} -> {} -> {} hitMask={} links=({}, {})", cellId, first.fromLayer, first.toLayer, second.toLayer, c.hitLayerMask.asString(), c.firstLink, c.secondLink); } return out; } @@ -113,15 +113,15 @@ class TrackingTopology for (int fromLayer = 0; fromLayer < mMaxLayers; ++fromLayer) { for (int toLayer = fromLayer + 1; toLayer < mMaxLayers; ++toLayer) { if (Mask::skipped(fromLayer, toLayer).isAllowedHoleMask(mMaxHoles, mHoleLayerMask)) { - mTransitions[mNTransitions++] = LayerTransition{static_cast(fromLayer), static_cast(toLayer)}; + mLinks[mNLinks++] = LayerLink{static_cast(fromLayer), static_cast(toLayer)}; } } } - for (Id firstId = 0; firstId < mNTransitions; ++firstId) { - const auto& first = mTransitions[firstId]; - for (Id secondId = 0; secondId < mNTransitions; ++secondId) { - const auto& second = mTransitions[secondId]; + for (Id firstId = 0; firstId < mNLinks; ++firstId) { + const auto& first = mLinks[firstId]; + for (Id secondId = 0; secondId < mNLinks; ++secondId) { + const auto& second = mLinks[secondId]; if (first.toLayer != second.fromLayer) { continue; } @@ -132,86 +132,86 @@ class TrackingTopology } } - fillCellsByTransition(); + fillCellsByLink(); } View getView() const { - return View{mTransitions.data(), + return View{mLinks.data(), mCells.data(), - mCellsByFirstTransitionIndex.data(), - mCellsByFirstTransition.data(), - mNTransitions, + mCellsByFirstLinkIndex.data(), + mCellsByFirstLink.data(), + mNLinks, mNCells, - mNCellsByFirstTransition}; + mNCellsByFirstLink}; } - View getDeviceView(const LayerTransition* deviceTransitions, + View getDeviceView(const LayerLink* deviceLinks, const CellTopology* deviceCells, - const Range* deviceCellsByFirstTransitionIndex, - const Id* deviceCellsByFirstTransition) const + const Range* deviceCellsByFirstLinkIndex, + const Id* deviceCellsByFirstLink) const { - return View{deviceTransitions, + return View{deviceLinks, deviceCells, - deviceCellsByFirstTransitionIndex, - deviceCellsByFirstTransition, - mNTransitions, + deviceCellsByFirstLinkIndex, + deviceCellsByFirstLink, + mNLinks, mNCells, - mNCellsByFirstTransition}; + mNCellsByFirstLink}; } - const auto& getTransitions() const noexcept { return mTransitions; } + const auto& getLinks() const noexcept { return mLinks; } const auto& getCells() const noexcept { return mCells; } - const auto& getCellsByFirstTransitionIndex() const noexcept { return mCellsByFirstTransitionIndex; } - const auto& getCellsByFirstTransition() const noexcept { return mCellsByFirstTransition; } - Id getNTransitions() const noexcept { return mNTransitions; } + const auto& getCellsByFirstLinkIndex() const noexcept { return mCellsByFirstLinkIndex; } + const auto& getCellsByFirstLink() const noexcept { return mCellsByFirstLink; } + Id getNLinks() const noexcept { return mNLinks; } Id getNCells() const noexcept { return mNCells; } - Id getNCellsByFirstTransition() const noexcept { return mNCellsByFirstTransition; } + Id getNCellsByFirstLink() const noexcept { return mNCellsByFirstLink; } private: void clear() { - mNTransitions = 0; + mNLinks = 0; mNCells = 0; - mNCellsByFirstTransition = 0; - mTransitions.fill({}); + mNCellsByFirstLink = 0; + mLinks.fill({}); mCells.fill({}); - mCellsByFirstTransitionIndex.fill(Range{0, 0}); - mCellsByFirstTransition.fill(0); + mCellsByFirstLinkIndex.fill(Range{0, 0}); + mCellsByFirstLink.fill(0); } - void fillCellsByTransition() + void fillCellsByLink() { - std::array counts{}; + std::array counts{}; for (Id cellId = 0; cellId < mNCells; ++cellId) { - ++counts[mCells[cellId].firstTransition]; + ++counts[mCells[cellId].firstLink]; } Id offset = 0; - for (Id transitionId = 0; transitionId < mNTransitions; ++transitionId) { - mCellsByFirstTransitionIndex[transitionId].setFirstEntry(offset); - mCellsByFirstTransitionIndex[transitionId].setEntries(counts[transitionId]); - offset += counts[transitionId]; + for (Id linkId = 0; linkId < mNLinks; ++linkId) { + mCellsByFirstLinkIndex[linkId].setFirstEntry(offset); + mCellsByFirstLinkIndex[linkId].setEntries(counts[linkId]); + offset += counts[linkId]; } - std::array cursor{}; + std::array cursor{}; for (Id cellId = 0; cellId < mNCells; ++cellId) { - const Id transitionId = mCells[cellId].firstTransition; - mCellsByFirstTransition[mCellsByFirstTransitionIndex[transitionId].getFirstEntry() + cursor[transitionId]++] = cellId; + const Id linkId = mCells[cellId].firstLink; + mCellsByFirstLink[mCellsByFirstLinkIndex[linkId].getFirstEntry() + cursor[linkId]++] = cellId; } - mNCellsByFirstTransition = offset; + mNCellsByFirstLink = offset; } int mMaxLayers{0}; int mMaxHoles{0}; Mask mHoleLayerMask{0}; - Id mNTransitions{0}; + Id mNLinks{0}; Id mNCells{0}; - Id mNCellsByFirstTransition{0}; - std::array mTransitions{}; + Id mNCellsByFirstLink{0}; + std::array mLinks{}; std::array mCells{}; - std::array mCellsByFirstTransitionIndex{}; - std::array mCellsByFirstTransition{}; + std::array mCellsByFirstLinkIndex{}; + std::array mCellsByFirstLink{}; }; } // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index 0442fec1a7631..d418911202991 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -312,11 +312,11 @@ void TimeFrame::initialise(const TrackingParameters& trkParam, const in clearResizeBoundedVector(mCellsNeighboursTopology, mTrackingTopologyView.nCells, mMemoryPool.get()); clearResizeBoundedVector(mCellsNeighboursLUT, mTrackingTopologyView.nCells, mMemoryPool.get()); clearResizeBoundedVector(mCellLabels, mTrackingTopologyView.nCells, mMemoryPool.get()); - clearResizeBoundedVector(mTracklets, mTrackingTopologyView.nTransitions, mMemoryPool.get()); - clearResizeBoundedVector(mTrackletLabels, mTrackingTopologyView.nTransitions, mMemoryPool.get()); - clearResizeBoundedVector(mTrackletsLookupTable, mTrackingTopologyView.nTransitions, mMemoryPool.get()); - clearResizeBoundedVector(mTransitionPhiCuts, mTrackingTopologyView.nTransitions, mMemoryPool.get()); - clearResizeBoundedVector(mTransitionMSAngles, mTrackingTopologyView.nTransitions, mMemoryPool.get()); + clearResizeBoundedVector(mTracklets, mTrackingTopologyView.nLinks, mMemoryPool.get()); + clearResizeBoundedVector(mTrackletLabels, mTrackingTopologyView.nLinks, mMemoryPool.get()); + clearResizeBoundedVector(mTrackletsLookupTable, mTrackingTopologyView.nLinks, mMemoryPool.get()); + clearResizeBoundedVector(mLinkPhiCuts, mTrackingTopologyView.nLinks, mMemoryPool.get()); + clearResizeBoundedVector(mLinkMSAngles, mTrackingTopologyView.nLinks, mMemoryPool.get()); mNTrackletsPerROF.resize(2); for (auto& v : mNTrackletsPerROF) { v = bounded_vector(getNrof(1) + 1, 0, mMemoryPool.get()); @@ -339,32 +339,32 @@ void TimeFrame::initialise(const TrackingParameters& trkParam, const in mPositionResolution[iLayer] = o2::gpu::CAMath::Sqrt((0.5f * (trkParam.SystErrorZ2[iLayer] + trkParam.SystErrorY2[iLayer])) + (trkParam.LayerResolution[iLayer] * trkParam.LayerResolution[iLayer])); } - // for each transition calculate the phi-cuts + integrated MS + // for each link calculate the phi-cuts + integrated MS float oneOverR{0.001f * 0.3f * std::abs(mBz) / trkParam.TrackletMinPt}; - for (int transitionId{0}; transitionId < (int)mTracklets.size(); ++transitionId) { - const auto& transition = mTrackingTopologyView.getTransition(transitionId); + for (int linkId{0}; linkId < (int)mTracklets.size(); ++linkId) { + const auto& link = mTrackingTopologyView.getLink(linkId); float ms2 = 0.; - for (int layer = transition.fromLayer; layer < transition.toLayer; ++layer) { + for (int layer = link.fromLayer; layer < link.toLayer; ++layer) { ms2 += math_utils::Sq(msAngles[layer]); } - mTransitionMSAngles[transitionId] = o2::gpu::CAMath::Sqrt(ms2); - const float& r1 = trkParam.LayerRadii[transition.fromLayer]; - const float& r2 = trkParam.LayerRadii[transition.toLayer]; + mLinkMSAngles[linkId] = o2::gpu::CAMath::Sqrt(ms2); + const float& r1 = trkParam.LayerRadii[link.fromLayer]; + const float& r2 = trkParam.LayerRadii[link.toLayer]; oneOverR = (0.5 * oneOverR >= 1.f / r2) ? (2.f / r2) - o2::constants::math::Almost0 : oneOverR; - const float res1 = o2::gpu::CAMath::Hypot(trkParam.PVres, mPositionResolution[transition.fromLayer]); - const float res2 = o2::gpu::CAMath::Hypot(trkParam.PVres, mPositionResolution[transition.toLayer]); + const float res1 = o2::gpu::CAMath::Hypot(trkParam.PVres, mPositionResolution[link.fromLayer]); + const float res2 = o2::gpu::CAMath::Hypot(trkParam.PVres, mPositionResolution[link.toLayer]); const float cosTheta1half = o2::gpu::CAMath::Sqrt(1.f - math_utils::Sq(0.5f * r1 * oneOverR)); const float cosTheta2half = o2::gpu::CAMath::Sqrt(1.f - math_utils::Sq(0.5f * r2 * oneOverR)); float x = (r2 * cosTheta1half) - (r1 * cosTheta2half); float delta = o2::gpu::CAMath::Sqrt(1.f / (1.f - 0.25f * math_utils::Sq(x * oneOverR)) * (math_utils::Sq((0.25f * r1 * r2 * math_utils::Sq(oneOverR) / cosTheta2half) + cosTheta1half) * math_utils::Sq(res1) + math_utils::Sq((0.25f * r1 * r2 * math_utils::Sq(oneOverR) / cosTheta1half) + cosTheta2half) * math_utils::Sq(res2))); /// the expression std::asin(0.5f * x * oneOverR) is equivalent to std::aCos(0.5f * r1 * oneOverR) - std::acos(0.5 * r2 * oneOverR) - mTransitionPhiCuts[transitionId] = o2::gpu::CAMath::Min(o2::gpu::CAMath::ASin(0.5f * x * oneOverR) + 2.f * mTransitionMSAngles[transitionId] + delta, o2::constants::math::PI * 0.5f); + mLinkPhiCuts[linkId] = o2::gpu::CAMath::Min(o2::gpu::CAMath::ASin(0.5f * x * oneOverR) + 2.f * mLinkMSAngles[linkId] + delta, o2::constants::math::PI * 0.5f); // some cleanup - deepVectorClear(mTracklets[transitionId]); - deepVectorClear(mTrackletLabels[transitionId]); - deepVectorClear(mTrackletsLookupTable[transitionId]); - mTrackletsLookupTable[transitionId].resize(mClusters[transition.fromLayer].size() + 1, 0); + deepVectorClear(mTracklets[linkId]); + deepVectorClear(mTrackletLabels[linkId]); + deepVectorClear(mTrackletsLookupTable[linkId]); + mTrackletsLookupTable[linkId].resize(mClusters[link.fromLayer].size() + 1, 0); } for (int cellId{0}; cellId < (int)mCells.size(); ++cellId) { @@ -438,8 +438,8 @@ void TimeFrame::setMemoryPool(std::shared_ptr po initContainers(mNTrackletsPerClusterSum); initContainers(mNClustersPerROF); initVector(mPrimaryVertices); - initVector(mTransitionPhiCuts); - initVector(mTransitionMSAngles); + initVector(mLinkPhiCuts); + initVector(mLinkMSAngles); initVector(mPositionResolution); initContainers(mClusterSize); initVector(mPValphaX); @@ -488,8 +488,8 @@ void TimeFrame::wipe() deepVectorClear(mNTrackletsPerCluster); deepVectorClear(mNTrackletsPerClusterSum); deepVectorClear(mNClustersPerROF); - deepVectorClear(mTransitionPhiCuts); - deepVectorClear(mTransitionMSAngles); + deepVectorClear(mLinkPhiCuts); + deepVectorClear(mLinkMSAngles); deepVectorClear(mPositionResolution); deepVectorClear(mClusterSize); deepVectorClear(mPValphaX); diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 2b5a17290b94a..e4523b903f075 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -52,22 +52,22 @@ template void TrackerTraits::computeLayerTracklets(const int iteration, int iVertex) { const auto topology = mTimeFrame->getTrackingTopologyView(); - for (int transitionId = 0; transitionId < topology.nTransitions; ++transitionId) { - mTimeFrame->getTracklets()[transitionId].clear(); - mTimeFrame->getTrackletsLabel(transitionId).clear(); - std::fill(mTimeFrame->getTrackletsLookupTable()[transitionId].begin(), mTimeFrame->getTrackletsLookupTable()[transitionId].end(), 0); + for (int linkId = 0; linkId < topology.nLinks; ++linkId) { + mTimeFrame->getTracklets()[linkId].clear(); + mTimeFrame->getTrackletsLabel(linkId).clear(); + std::fill(mTimeFrame->getTrackletsLookupTable()[linkId].begin(), mTimeFrame->getTrackletsLookupTable()[linkId].end(), 0); } const Vertex diamondVert(mTrkParams[iteration].Diamond, mTrkParams[iteration].DiamondCov, 1, 1.f); gsl::span diamondSpan(&diamondVert, 1); mTaskArena->execute([&] { - auto forTracklets = [&](auto Tag, int transitionId, int pivotROF, int base, int& offset) -> int { - const auto& transition = topology.getTransition(transitionId); - if (!mTimeFrame->getROFMaskView().isROFEnabled(transition.fromLayer, pivotROF)) { + auto forTracklets = [&](auto Tag, int linkId, int pivotROF, int base, int& offset) -> int { + const auto& link = topology.getLink(linkId); + if (!mTimeFrame->getROFMaskView().isROFEnabled(link.fromLayer, pivotROF)) { return 0; } - gsl::span primaryVertices = mTrkParams[iteration].UseDiamond ? diamondSpan : mTimeFrame->getPrimaryVertices(transition.fromLayer, pivotROF); + gsl::span primaryVertices = mTrkParams[iteration].UseDiamond ? diamondSpan : mTimeFrame->getPrimaryVertices(link.fromLayer, pivotROF); if (primaryVertices.empty()) { return 0; } @@ -77,45 +77,45 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer return 0; } - const auto& rofOverlap = mTimeFrame->getROFOverlapTableView().getOverlap(transition.fromLayer, transition.toLayer, pivotROF); + const auto& rofOverlap = mTimeFrame->getROFOverlapTableView().getOverlap(link.fromLayer, link.toLayer, pivotROF); if (!rofOverlap.getEntries()) { return 0; } int localCount = 0; - auto& tracklets = mTimeFrame->getTracklets()[transitionId]; - auto layer0 = mTimeFrame->getClustersOnLayer(pivotROF, transition.fromLayer); + auto& tracklets = mTimeFrame->getTracklets()[linkId]; + auto layer0 = mTimeFrame->getClustersOnLayer(pivotROF, link.fromLayer); if (layer0.empty()) { return 0; } - const float meanDeltaR = mTrkParams[iteration].LayerRadii[transition.toLayer] - mTrkParams[iteration].LayerRadii[transition.fromLayer]; - const float phiCut = mTimeFrame->getTransitionPhiCut(transitionId); - const float msAngle = mTimeFrame->getTransitionMSAngle(transitionId); + const float meanDeltaR = mTrkParams[iteration].LayerRadii[link.toLayer] - mTrkParams[iteration].LayerRadii[link.fromLayer]; + const float phiCut = mTimeFrame->getLinkPhiCut(linkId); + const float msAngle = mTimeFrame->getLinkMSAngle(linkId); for (int iCluster = 0; iCluster < int(layer0.size()); ++iCluster) { const Cluster& currentCluster = layer0[iCluster]; - const int currentSortedIndex = mTimeFrame->getSortedIndex(pivotROF, transition.fromLayer, iCluster); - if (mTimeFrame->isClusterUsed(transition.fromLayer, currentCluster.clusterId)) { + const int currentSortedIndex = mTimeFrame->getSortedIndex(pivotROF, link.fromLayer, iCluster); + if (mTimeFrame->isClusterUsed(link.fromLayer, currentCluster.clusterId)) { continue; } const float inverseR0 = 1.f / currentCluster.radius; for (int iV = startVtx; iV < endVtx; ++iV) { const auto& pv = primaryVertices[iV]; - if (!mTimeFrame->getROFVertexLookupTableView().isVertexCompatible(transition.fromLayer, pivotROF, pv)) { + if (!mTimeFrame->getROFVertexLookupTableView().isVertexCompatible(link.fromLayer, pivotROF, pv)) { continue; } if (pv.isFlagSet(Vertex::Flags::UPCMode) != mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices]) { continue; } - const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(mTimeFrame->getPositionResolution(transition.fromLayer)) + math_utils::Sq(mTrkParams[iteration].PVres) / float(pv.getNContributors())); + const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(mTimeFrame->getPositionResolution(link.fromLayer)) + math_utils::Sq(mTrkParams[iteration].PVres) / float(pv.getNContributors())); const float tanLambda = (currentCluster.zCoordinate - pv.getZ()) * inverseR0; - const float zAtRmin = tanLambda * (mTimeFrame->getMinR(transition.toLayer) - currentCluster.radius) + currentCluster.zCoordinate; - const float zAtRmax = tanLambda * (mTimeFrame->getMaxR(transition.toLayer) - currentCluster.radius) + currentCluster.zCoordinate; + const float zAtRmin = tanLambda * (mTimeFrame->getMinR(link.toLayer) - currentCluster.radius) + currentCluster.zCoordinate; + const float zAtRmax = tanLambda * (mTimeFrame->getMaxR(link.toLayer) - currentCluster.radius) + currentCluster.zCoordinate; const float sqInvDeltaZ0 = 1.f / (math_utils::Sq(currentCluster.zCoordinate - pv.getZ()) + constants::Tolerance); const float sigmaZ = o2::gpu::CAMath::Sqrt((math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInvDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f)) + math_utils::Sq(meanDeltaR * msAngle)); - const auto bins = o2::its::getBinsRect(currentCluster, transition.toLayer, zAtRmin, zAtRmax, + const auto bins = o2::its::getBinsRect(currentCluster, link.toLayer, zAtRmin, zAtRmax, sigmaZ * mTrkParams[iteration].NSigmaCut, phiCut, mTimeFrame->getIndexTableUtils()); if (bins.x < 0) { @@ -127,18 +127,18 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer } for (int targetROF = rofOverlap.getFirstEntry(); targetROF < rofOverlap.getEntriesBound(); ++targetROF) { - if (!mTimeFrame->getROFMaskView().isROFEnabled(transition.toLayer, targetROF)) { + if (!mTimeFrame->getROFMaskView().isROFEnabled(link.toLayer, targetROF)) { continue; } - auto layer1 = mTimeFrame->getClustersOnLayer(targetROF, transition.toLayer); + auto layer1 = mTimeFrame->getClustersOnLayer(targetROF, link.toLayer); if (layer1.empty()) { continue; } - const auto ts = mTimeFrame->getROFOverlapTableView().getTimeStamp(transition.fromLayer, pivotROF, transition.toLayer, targetROF); + const auto ts = mTimeFrame->getROFOverlapTableView().getTimeStamp(link.fromLayer, pivotROF, link.toLayer, targetROF); if (!ts.isCompatible(pv.getTimeStamp())) { continue; } - const auto& targetIndexTable = mTimeFrame->getIndexTable(targetROF, transition.toLayer); + const auto& targetIndexTable = mTimeFrame->getIndexTable(targetROF, link.toLayer); const int zBinRange = (bins.z - bins.x) + 1; for (int iPhi = 0; iPhi < phiBinsNum; ++iPhi) { const int iPhiBin = (bins.y + iPhi) % mTrkParams[iteration].PhiBins; @@ -151,7 +151,7 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer break; } const Cluster& nextCluster = layer1[iNext]; - if (mTimeFrame->isClusterUsed(transition.toLayer, nextCluster.clusterId)) { + if (mTimeFrame->isClusterUsed(link.toLayer, nextCluster.clusterId)) { continue; } const float deltaZ = o2::gpu::CAMath::Abs((tanLambda * (nextCluster.radius - currentCluster.radius)) + currentCluster.zCoordinate - nextCluster.zCoordinate); @@ -161,12 +161,12 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer const float phi{o2::gpu::CAMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate)}; const float tanL = (currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius); if constexpr (decltype(Tag)::value == PassMode::OnePass::value) { - tracklets.emplace_back(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF, transition.toLayer, iNext), tanL, phi, ts); + tracklets.emplace_back(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF, link.toLayer, iNext), tanL, phi, ts); } else if constexpr (decltype(Tag)::value == PassMode::TwoPassCount::value) { ++localCount; } else if constexpr (decltype(Tag)::value == PassMode::TwoPassInsert::value) { const int idx = base + offset++; - tracklets[idx] = Tracklet(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF, transition.toLayer, iNext), tanL, phi, ts); + tracklets[idx] = Tracklet(currentSortedIndex, mTimeFrame->getSortedIndex(targetROF, link.toLayer, iNext), tanL, phi, ts); } } } @@ -179,24 +179,24 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer int dummy{0}; if (mTaskArena->max_concurrency() <= 1) { - for (int transitionId{0}; transitionId < topology.nTransitions; ++transitionId) { - const int fromLayer = topology.getTransition(transitionId).fromLayer; + for (int linkId{0}; linkId < topology.nLinks; ++linkId) { + const int fromLayer = topology.getLink(linkId).fromLayer; const int startROF = 0, endROF = mTimeFrame->getROFOverlapTableView().getLayer(fromLayer).mNROFsTF; for (int pivotROF{startROF}; pivotROF < endROF; ++pivotROF) { - forTracklets(PassMode::OnePass{}, transitionId, pivotROF, 0, dummy); + forTracklets(PassMode::OnePass{}, linkId, pivotROF, 0, dummy); } } } else { - tbb::parallel_for(0, static_cast(topology.nTransitions), [&](const int transitionId) { - const int fromLayer = topology.getTransition(transitionId).fromLayer; + tbb::parallel_for(0, static_cast(topology.nLinks), [&](const int linkId) { + const int fromLayer = topology.getLink(linkId).fromLayer; const int startROF = 0, endROF = mTimeFrame->getROFOverlapTableView().getLayer(fromLayer).mNROFsTF; bounded_vector perROFCount((endROF - startROF) + 1, mMemoryPool.get()); tbb::parallel_for(startROF, endROF, [&](const int pivotROF) { - perROFCount[pivotROF - startROF] = forTracklets(PassMode::TwoPassCount{}, transitionId, pivotROF, 0, dummy); + perROFCount[pivotROF - startROF] = forTracklets(PassMode::TwoPassCount{}, linkId, pivotROF, 0, dummy); }); std::exclusive_scan(perROFCount.begin(), perROFCount.end(), perROFCount.begin(), 0); const int nTracklets = perROFCount.back(); - mTimeFrame->getTracklets()[transitionId].resize(nTracklets); + mTimeFrame->getTracklets()[linkId].resize(nTracklets); if (nTracklets == 0) { return; } @@ -206,19 +206,19 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer return; } int localIdx = 0; - forTracklets(PassMode::TwoPassInsert{}, transitionId, pivotROF, baseIdx, localIdx); + forTracklets(PassMode::TwoPassInsert{}, linkId, pivotROF, baseIdx, localIdx); }); }); } - tbb::parallel_for(0, static_cast(topology.nTransitions), [&](const int transitionId) { + tbb::parallel_for(0, static_cast(topology.nLinks), [&](const int linkId) { /// Sort tracklets & remove duplicates // duplicates can exist simply since we evaluate per vertex - auto& trkl{mTimeFrame->getTracklets()[transitionId]}; + auto& trkl{mTimeFrame->getTracklets()[linkId]}; std::sort(trkl.begin(), trkl.end()); trkl.erase(std::unique(trkl.begin(), trkl.end()), trkl.end()); trkl.shrink_to_fit(); - auto& lut{mTimeFrame->getTrackletsLookupTable()[transitionId]}; + auto& lut{mTimeFrame->getTrackletsLookupTable()[linkId]}; if (!trkl.empty()) { for (const auto& tkl : trkl) { lut[tkl.firstClusterIndex + 1]++; @@ -229,14 +229,14 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer /// Create tracklets labels if (mTimeFrame->hasMCinformation() && mTrkParams[iteration].CreateArtefactLabels) { - tbb::parallel_for(0, static_cast(topology.nTransitions), [&](const int transitionId) { - const auto& transition = topology.getTransition(transitionId); - for (auto& trk : mTimeFrame->getTracklets()[transitionId]) { + tbb::parallel_for(0, static_cast(topology.nLinks), [&](const int linkId) { + const auto& link = topology.getLink(linkId); + for (auto& trk : mTimeFrame->getTracklets()[linkId]) { MCCompLabel label; - int currentId{mTimeFrame->getClusters()[transition.fromLayer][trk.firstClusterIndex].clusterId}; - int nextId{mTimeFrame->getClusters()[transition.toLayer][trk.secondClusterIndex].clusterId}; - for (const auto& lab1 : mTimeFrame->getClusterLabels(transition.fromLayer, currentId)) { - for (const auto& lab2 : mTimeFrame->getClusterLabels(transition.toLayer, nextId)) { + int currentId{mTimeFrame->getClusters()[link.fromLayer][trk.firstClusterIndex].clusterId}; + int nextId{mTimeFrame->getClusters()[link.toLayer][trk.secondClusterIndex].clusterId}; + for (const auto& lab1 : mTimeFrame->getClusterLabels(link.fromLayer, currentId)) { + for (const auto& lab2 : mTimeFrame->getClusterLabels(link.toLayer, nextId)) { if (lab1 == lab2 && lab1.isValid()) { label = lab1; break; @@ -246,7 +246,7 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iVer break; } } - mTimeFrame->getTrackletsLabel(transitionId).emplace_back(label); + mTimeFrame->getTrackletsLabel(linkId).emplace_back(label); } }); } @@ -268,15 +268,15 @@ void TrackerTraits::computeLayerCells(const int iteration) mTaskArena->execute([&] { auto forTrackletCells = [&](auto Tag, int cellTopologyId, bounded_vector& layerCells, int iTracklet, int offset = 0) -> int { const auto& cellTopology = topology.getCell(cellTopologyId); - const auto& firstTransition = topology.getTransition(cellTopology.firstTransition); - const auto& secondTransition = topology.getTransition(cellTopology.secondTransition); - const Tracklet& currentTracklet{mTimeFrame->getTracklets()[cellTopology.firstTransition][iTracklet]}; + const auto& firstLink = topology.getLink(cellTopology.firstLink); + const auto& secondLink = topology.getLink(cellTopology.secondLink); + const Tracklet& currentTracklet{mTimeFrame->getTracklets()[cellTopology.firstLink][iTracklet]}; const int nextLayerClusterIndex{currentTracklet.secondClusterIndex}; - const int nextLayerFirstTrackletIndex{mTimeFrame->getTrackletsLookupTable()[cellTopology.secondTransition][nextLayerClusterIndex]}; - const int nextLayerLastTrackletIndex{mTimeFrame->getTrackletsLookupTable()[cellTopology.secondTransition][nextLayerClusterIndex + 1]}; + const int nextLayerFirstTrackletIndex{mTimeFrame->getTrackletsLookupTable()[cellTopology.secondLink][nextLayerClusterIndex]}; + const int nextLayerLastTrackletIndex{mTimeFrame->getTrackletsLookupTable()[cellTopology.secondLink][nextLayerClusterIndex + 1]}; int foundCells{0}; for (int iNextTracklet{nextLayerFirstTrackletIndex}; iNextTracklet < nextLayerLastTrackletIndex; ++iNextTracklet) { - const Tracklet& nextTracklet{mTimeFrame->getTracklets()[cellTopology.secondTransition][iNextTracklet]}; + const Tracklet& nextTracklet{mTimeFrame->getTracklets()[cellTopology.secondLink][iNextTracklet]}; if (nextTracklet.firstClusterIndex != nextLayerClusterIndex) { break; } @@ -289,13 +289,13 @@ void TrackerTraits::computeLayerCells(const int iteration) /// Track seed preparation. Clusters are numbered progressively from the innermost going outward. const int clusId[3]{ - mTimeFrame->getClusters()[firstTransition.fromLayer][currentTracklet.firstClusterIndex].clusterId, - mTimeFrame->getClusters()[firstTransition.toLayer][nextTracklet.firstClusterIndex].clusterId, - mTimeFrame->getClusters()[secondTransition.toLayer][nextTracklet.secondClusterIndex].clusterId}; - const int hitLayers[3]{firstTransition.fromLayer, firstTransition.toLayer, secondTransition.toLayer}; - const auto& cluster1_glo = mTimeFrame->getUnsortedClusters()[firstTransition.fromLayer][clusId[0]]; - const auto& cluster2_glo = mTimeFrame->getUnsortedClusters()[firstTransition.toLayer][clusId[1]]; - const auto& cluster3_tf = mTimeFrame->getTrackingFrameInfoOnLayer(secondTransition.toLayer)[clusId[2]]; + mTimeFrame->getClusters()[firstLink.fromLayer][currentTracklet.firstClusterIndex].clusterId, + mTimeFrame->getClusters()[firstLink.toLayer][nextTracklet.firstClusterIndex].clusterId, + mTimeFrame->getClusters()[secondLink.toLayer][nextTracklet.secondClusterIndex].clusterId}; + const int hitLayers[3]{firstLink.fromLayer, firstLink.toLayer, secondLink.toLayer}; + const auto& cluster1_glo = mTimeFrame->getUnsortedClusters()[firstLink.fromLayer][clusId[0]]; + const auto& cluster2_glo = mTimeFrame->getUnsortedClusters()[firstLink.toLayer][clusId[1]]; + const auto& cluster3_tf = mTimeFrame->getTrackingFrameInfoOnLayer(secondLink.toLayer)[clusId[2]]; auto track{o2::its::track::buildTrackSeed(cluster1_glo, cluster2_glo, cluster3_tf, mBz)}; float chi2{0.f}; @@ -350,13 +350,13 @@ void TrackerTraits::computeLayerCells(const int iteration) for (int cellTopologyId = 0; cellTopologyId < topology.nCells; ++cellTopologyId) { const auto& cellTopology = topology.getCell(cellTopologyId); - if (mTimeFrame->getTracklets()[cellTopology.firstTransition].empty() || - mTimeFrame->getTracklets()[cellTopology.secondTransition].empty()) { + if (mTimeFrame->getTracklets()[cellTopology.firstLink].empty() || + mTimeFrame->getTracklets()[cellTopology.secondLink].empty()) { continue; } auto& layerCells = mTimeFrame->getCells()[cellTopologyId]; - const int currentLayerTrackletsNum{static_cast(mTimeFrame->getTracklets()[cellTopology.firstTransition].size())}; + const int currentLayerTrackletsNum{static_cast(mTimeFrame->getTracklets()[cellTopology.firstLink].size())}; bounded_vector perTrackletCount(currentLayerTrackletsNum + 1, 0, mMemoryPool.get()); if (mTaskArena->max_concurrency() <= 1) { for (int iTracklet{0}; iTracklet < currentLayerTrackletsNum; ++iTracklet) { @@ -395,17 +395,17 @@ void TrackerTraits::computeLayerCells(const int iteration) auto& labels = mTimeFrame->getCellsLabel(cellTopologyId); labels.reserve(layerCells.size()); for (const auto& cell : layerCells) { - MCCompLabel currentLab{mTimeFrame->getTrackletsLabel(cellTopology.firstTransition)[cell.getFirstTrackletIndex()]}; - MCCompLabel nextLab{mTimeFrame->getTrackletsLabel(cellTopology.secondTransition)[cell.getSecondTrackletIndex()]}; + MCCompLabel currentLab{mTimeFrame->getTrackletsLabel(cellTopology.firstLink)[cell.getFirstTrackletIndex()]}; + MCCompLabel nextLab{mTimeFrame->getTrackletsLabel(cellTopology.secondLink)[cell.getSecondTrackletIndex()]}; labels.emplace_back(currentLab == nextLab ? currentLab : MCCompLabel()); } } } }); - for (int transitionId = 0; transitionId < topology.nTransitions; ++transitionId) { - deepVectorClear(mTimeFrame->getTracklets()[transitionId]); - deepVectorClear(mTimeFrame->getTrackletsLabel(transitionId)); + for (int linkId = 0; linkId < topology.nLinks; ++linkId) { + deepVectorClear(mTimeFrame->getTracklets()[linkId]); + deepVectorClear(mTimeFrame->getTrackletsLabel(linkId)); } } @@ -430,7 +430,7 @@ void TrackerTraits::findCellsNeighbours(const int iteration) mTimeFrame->getCells()[cellTopologyId].empty()) { continue; } - const auto successors = topology.getCellsStartingWithTransition(cellTopology.secondTransition); + const auto successors = topology.getCellsStartingWithLink(cellTopology.secondLink); if (!successors.getEntries()) { continue; } @@ -441,7 +441,7 @@ void TrackerTraits::findCellsNeighbours(const int iteration) const auto& currentCellSeed{mTimeFrame->getCells()[cellTopologyId][iCell]}; const int nextLayerTrackletIndex{currentCellSeed.getSecondTrackletIndex()}; for (int iSuccessor{0}; iSuccessor < successors.getEntries(); ++iSuccessor) { - const int nextCellTopologyId = topology.cellsByFirstTransition[successors.getFirstEntry() + iSuccessor]; + const int nextCellTopologyId = topology.cellsByFirstLink[successors.getFirstEntry() + iSuccessor]; if (mTimeFrame->getCells()[nextCellTopologyId].empty() || mTimeFrame->getCellsLookupTable()[nextCellTopologyId].empty()) { continue; diff --git a/Detectors/ITSMFT/ITS/tracking/test/testTrackingTopology.cxx b/Detectors/ITSMFT/ITS/tracking/test/testTrackingTopology.cxx index 4944d00b15fea..d3f249650a287 100644 --- a/Detectors/ITSMFT/ITS/tracking/test/testTrackingTopology.cxx +++ b/Detectors/ITSMFT/ITS/tracking/test/testTrackingTopology.cxx @@ -57,9 +57,9 @@ BOOST_AUTO_TEST_CASE(trackingtopology_basic) const auto view = topo.getView(); view.print(); - BOOST_CHECK_EQUAL(view.nTransitions, 3); + BOOST_CHECK_EQUAL(view.nLinks, 3); for (int i{0}; i < 3; ++i) { - const auto& tra = view.getTransition(i); + const auto& tra = view.getLink(i); BOOST_CHECK_EQUAL(tra.fromLayer, i); BOOST_CHECK_EQUAL(tra.toLayer, i + 1); } @@ -67,8 +67,8 @@ BOOST_AUTO_TEST_CASE(trackingtopology_basic) BOOST_CHECK_EQUAL(view.nCells, 2); for (int i{0}; i < 2; ++i) { const auto& cell = view.getCell(i); - BOOST_CHECK_EQUAL(cell.firstTransition, i); - BOOST_CHECK_EQUAL(cell.secondTransition, i + 1); + BOOST_CHECK_EQUAL(cell.firstLink, i); + BOOST_CHECK_EQUAL(cell.secondLink, i + 1); } } @@ -79,16 +79,16 @@ BOOST_AUTO_TEST_CASE(trackingtopology_single_allowed_hole) const auto view = topo.getView(); view.print(); - BOOST_CHECK_EQUAL(view.nTransitions, 5); + BOOST_CHECK_EQUAL(view.nLinks, 5); BOOST_CHECK_EQUAL(view.nCells, 5); - bool hasHoleTransition = false; - for (int i{0}; i < view.nTransitions; ++i) { - const auto& transition = view.getTransition(i); - hasHoleTransition |= transition.fromLayer == 1 && transition.toLayer == 3; - BOOST_CHECK(o2::its::LayerMask::skipped(transition.fromLayer, transition.toLayer).isAllowedHoleMask(1, 1 << 2)); + bool hasHoleLink = false; + for (int i{0}; i < view.nLinks; ++i) { + const auto& link = view.getLink(i); + hasHoleLink |= link.fromLayer == 1 && link.toLayer == 3; + BOOST_CHECK(o2::its::LayerMask::skipped(link.fromLayer, link.toLayer).isAllowedHoleMask(1, 1 << 2)); } - BOOST_CHECK(hasHoleTransition); + BOOST_CHECK(hasHoleLink); bool hasHoleCell = false; for (int i{0}; i < view.nCells; ++i) { @@ -106,10 +106,10 @@ BOOST_AUTO_TEST_CASE(trackingtopology_rejects_wrong_hole_layer) const auto view = topo.getView(); view.print(); - for (int i{0}; i < view.nTransitions; ++i) { - const auto& transition = view.getTransition(i); - BOOST_CHECK(!(transition.fromLayer == 0 && transition.toLayer == 2)); - BOOST_CHECK(!(transition.fromLayer == 2 && transition.toLayer == 4)); + for (int i{0}; i < view.nLinks; ++i) { + const auto& link = view.getLink(i); + BOOST_CHECK(!(link.fromLayer == 0 && link.toLayer == 2)); + BOOST_CHECK(!(link.fromLayer == 2 && link.toLayer == 4)); } for (int i{0}; i < view.nCells; ++i) {