Skip to content

ITS: restore previous UPC iteration #15289

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
f3sch wants to merge 4 commits into
base: dev
Choose a base branch
from
Open

ITS: restore previous UPC iteration #15289

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
7e53089
ITS: fix upc iteration
f3sch Apr 16, 2026
96cc11f
ITS: simplify configurables to single one for Vertexer
f3sch Apr 17, 2026
1dfdd2f
ITS: make MaxIter a constant
f3sch Apr 17, 2026
dd8f6fe
ITS: remove GPU params
f3sch Apr 17, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 3 additions & 21 deletions Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,8 +60,6 @@ void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
std::vector<float>& radii,
bounded_vector<float>& mulScatAng,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads,
gpu::Streams& streams);

template <int NLayers = 7>
Expand Down Expand Up @@ -93,8 +91,6 @@ void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
std::vector<float>& radii,
bounded_vector<float>& mulScatAng,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads,
gpu::Streams& streams);

template <int NLayers>
Expand All @@ -113,8 +109,6 @@ void countCellsHandler(const Cluster** sortedClusters,
const float cellDeltaTanLambdaSigma,
const float nSigmaCut,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads,
gpu::Streams& streams);

template <int NLayers>
Expand All @@ -132,8 +126,6 @@ void computeCellsHandler(const Cluster** sortedClusters,
const float maxChi2ClusterAttachment,
const float cellDeltaTanLambdaSigma,
const float nSigmaCut,
const int nBlocks,
const int nThreads,
gpu::Streams& streams);

template <int NLayers>
Expand All @@ -150,8 +142,6 @@ void countCellNeighboursHandler(CellSeed<NLayers>** cellsLayersDevice,
const unsigned int nCellsNext,
const int maxCellNeighbours,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads,
gpu::Stream& stream);

template <int NLayers>
Expand All @@ -167,8 +157,6 @@ void computeCellNeighboursHandler(CellSeed<NLayers>** cellsLayersDevice,
const unsigned int nCells,
const unsigned int nCellsNext,
const int maxCellNeighbours,
const int nBlocks,
const int nThreads,
gpu::Stream& stream);

int filterCellNeighboursHandler(gpuPair<int, int>*,
Expand All @@ -193,9 +181,7 @@ void processNeighboursHandler(const int startLayer,
const float maxChi2NDF,
const o2::base::Propagator* propagator,
const o2::base::PropagatorF::MatCorrType matCorrType,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads);
o2::its::ExternalAllocator* alloc);

template <int NLayers = 7>
void countTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
Expand All @@ -214,9 +200,7 @@ void countTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
const bool shiftRefToCluster,
const o2::base::Propagator* propagator,
const o2::base::PropagatorF::MatCorrType matCorrType,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads);
o2::its::ExternalAllocator* alloc);

template <int NLayers = 7>
void computeTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
Expand All @@ -237,9 +221,7 @@ void computeTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
const bool shiftRefToCluster,
const o2::base::Propagator* propagator,
const o2::base::PropagatorF::MatCorrType matCorrType,
o2::its::ExternalAllocator* alloc,
const int nBlocks,
const int nThreads);
o2::its::ExternalAllocator* alloc);

} // namespace o2::its
#endif // ITSTRACKINGGPU_TRACKINGKERNELS_H_
3 changes: 1 addition & 2 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -627,8 +627,7 @@ constexpr auto makeIterTags(std::index_sequence<I...>)
{
return std::array<uint64_t, sizeof...(I)>{makeIterTag<I>()...};
}
// FIXME: we have to be careful that the MaxIter does not diverge from the 4 here!
constexpr auto kIterTags = makeIterTags(std::make_index_sequence<4>{});
constexpr auto kIterTags = makeIterTags(std::make_index_sequence<constants::MaxIter>{});
} // namespace detail

template <int NLayers>
Expand Down
32 changes: 3 additions & 29 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,6 @@

#include "ITStrackingGPU/TrackerTraitsGPU.h"
#include "ITStrackingGPU/TrackingKernels.h"
#include "ITStracking/TrackingConfigParam.h"
#include "ITStracking/Constants.h"

namespace o2::its
Expand Down Expand Up @@ -63,8 +62,6 @@ void TrackerTraitsGPU<NLayers>::adoptTimeFrame(TimeFrame<NLayers>* tf)
template <int NLayers>
void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int iVertex)
{
const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();

// start by queuing loading needed of two last layers
for (int iLayer{NLayers}; iLayer-- > NLayers - 2;) {
mTimeFrameGPU->createUsedClustersDevice(iteration, iLayer);
Expand Down Expand Up @@ -109,8 +106,6 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
this->mTrkParams[iteration].LayerRadii,
mTimeFrameGPU->getMSangles(),
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksLayerTracklets[iteration],
conf.nThreadsLayerTracklets[iteration],
mTimeFrameGPU->getStreams());
mTimeFrameGPU->createTrackletsBuffers(iLayer);
if (mTimeFrameGPU->getNTracklets()[iLayer] == 0) {
Expand Down Expand Up @@ -144,17 +139,13 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
this->mTrkParams[iteration].LayerRadii,
mTimeFrameGPU->getMSangles(),
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksLayerTracklets[iteration],
conf.nThreadsLayerTracklets[iteration],
mTimeFrameGPU->getStreams());
}
}

template <int NLayers>
void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
{
auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();

// start by queuing loading needed of three last layers
for (int iLayer{NLayers}; iLayer-- > NLayers - 3;) {
mTimeFrameGPU->loadUnsortedClustersDevice(iteration, iLayer);
Expand Down Expand Up @@ -194,8 +185,6 @@ void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
this->mTrkParams[iteration].NSigmaCut,
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksLayerCells[iteration],
conf.nThreadsLayerCells[iteration],
mTimeFrameGPU->getStreams());
mTimeFrameGPU->createCellsBuffers(iLayer);
if (mTimeFrameGPU->getNCells()[iLayer] == 0) {
Expand All @@ -215,17 +204,13 @@ void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
this->mTrkParams[iteration].NSigmaCut,
conf.nBlocksLayerCells[iteration],
conf.nThreadsLayerCells[iteration],
mTimeFrameGPU->getStreams());
}
}

template <int NLayers>
void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
{
const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();

for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NeighboursPerRoad(); ++iLayer) {
const int currentLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer])};
const int nextLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer + 1])};
Expand All @@ -248,8 +233,6 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
nextLayerCellsNum,
1e2,
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksFindNeighbours[iteration],
conf.nThreadsFindNeighbours[iteration],
mTimeFrameGPU->getStream(iLayer));
mTimeFrameGPU->createNeighboursDevice(iLayer);
if (mTimeFrameGPU->getNNeighbours()[iLayer] == 0) {
Expand All @@ -267,8 +250,6 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
currentLayerCellsNum,
nextLayerCellsNum,
1e2,
conf.nBlocksFindNeighbours[iteration],
conf.nThreadsFindNeighbours[iteration],
mTimeFrameGPU->getStream(iLayer));
mTimeFrameGPU->getArrayNNeighbours()[iLayer] = filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
mTimeFrameGPU->getDeviceNeighbours(iLayer),
Expand All @@ -282,7 +263,6 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
template <int NLayers>
void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
{
auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
const int minimumLayer{startLevel - 1};
bounded_vector<CellSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
Expand All @@ -305,9 +285,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
this->mTrkParams[0].MaxChi2NDF,
mTimeFrameGPU->getDevicePropagator(),
this->mTrkParams[0].CorrType,
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksProcessNeighbours[iteration],
conf.nThreadsProcessNeighbours[iteration]);
mTimeFrameGPU->getFrameworkAllocator());
}
// fixme: I don't want to move tracks back and forth, but I need a way to use a thrust::allocator that is aware of our managed memory.
if (trackSeeds.empty()) {
Expand All @@ -334,9 +312,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
this->mTrkParams[0].ShiftRefToCluster,
mTimeFrameGPU->getDevicePropagator(),
this->mTrkParams[0].CorrType,
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksTracksSeeds[iteration],
conf.nThreadsTracksSeeds[iteration]);
mTimeFrameGPU->getFrameworkAllocator());
mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size());
computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
Expand All @@ -356,9 +332,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
this->mTrkParams[0].ShiftRefToCluster,
mTimeFrameGPU->getDevicePropagator(),
this->mTrkParams[0].CorrType,
mTimeFrameGPU->getFrameworkAllocator(),
conf.nBlocksTracksSeeds[iteration],
conf.nThreadsTracksSeeds[iteration]);
mTimeFrameGPU->getFrameworkAllocator());
mTimeFrameGPU->downloadTrackITSExtDevice();

auto& tracks = mTimeFrameGPU->getTrackITSExt();
Expand Down
Loading