Skip to content

Commit ded827e

Browse files
f3schf3sch
authored
ITS: restore previous UPC iteration (#15289)
* ITS: fix upc iteration Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: simplify configurables to single one for Vertexer Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: make MaxIter a constant Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITS: remove GPU params Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> --------- Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent a105f1e commit ded827e

File tree

18 files changed

+118
-271
lines changed

18 files changed

+118
-271
lines changed

‎Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h‎

Lines changed: 3 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -60,8 +60,6 @@ void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
6060
std::vector<float>& radii,
6161
bounded_vector<float>& mulScatAng,
6262
o2::its::ExternalAllocator* alloc,
63-
const int nBlocks,
64-
const int nThreads,
6563
gpu::Streams& streams);
6664

6765
template <int NLayers = 7>
@@ -93,8 +91,6 @@ void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
9391
std::vector<float>& radii,
9492
bounded_vector<float>& mulScatAng,
9593
o2::its::ExternalAllocator* alloc,
96-
const int nBlocks,
97-
const int nThreads,
9894
gpu::Streams& streams);
9995

10096
template <int NLayers>
@@ -113,8 +109,6 @@ void countCellsHandler(const Cluster** sortedClusters,
113109
const float cellDeltaTanLambdaSigma,
114110
const float nSigmaCut,
115111
o2::its::ExternalAllocator* alloc,
116-
const int nBlocks,
117-
const int nThreads,
118112
gpu::Streams& streams);
119113

120114
template <int NLayers>
@@ -132,8 +126,6 @@ void computeCellsHandler(const Cluster** sortedClusters,
132126
const float maxChi2ClusterAttachment,
133127
const float cellDeltaTanLambdaSigma,
134128
const float nSigmaCut,
135-
const int nBlocks,
136-
const int nThreads,
137129
gpu::Streams& streams);
138130

139131
template <int NLayers>
@@ -150,8 +142,6 @@ void countCellNeighboursHandler(CellSeed<NLayers>** cellsLayersDevice,
150142
const unsigned int nCellsNext,
151143
const int maxCellNeighbours,
152144
o2::its::ExternalAllocator* alloc,
153-
const int nBlocks,
154-
const int nThreads,
155145
gpu::Stream& stream);
156146

157147
template <int NLayers>
@@ -167,8 +157,6 @@ void computeCellNeighboursHandler(CellSeed<NLayers>** cellsLayersDevice,
167157
const unsigned int nCells,
168158
const unsigned int nCellsNext,
169159
const int maxCellNeighbours,
170-
const int nBlocks,
171-
const int nThreads,
172160
gpu::Stream& stream);
173161

174162
int filterCellNeighboursHandler(gpuPair<int, int>*,
@@ -193,9 +181,7 @@ void processNeighboursHandler(const int startLayer,
193181
const float maxChi2NDF,
194182
const o2::base::Propagator* propagator,
195183
const o2::base::PropagatorF::MatCorrType matCorrType,
196-
o2::its::ExternalAllocator* alloc,
197-
const int nBlocks,
198-
const int nThreads);
184+
o2::its::ExternalAllocator* alloc);
199185

200186
template <int NLayers = 7>
201187
void countTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
@@ -214,9 +200,7 @@ void countTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
214200
const bool shiftRefToCluster,
215201
const o2::base::Propagator* propagator,
216202
const o2::base::PropagatorF::MatCorrType matCorrType,
217-
o2::its::ExternalAllocator* alloc,
218-
const int nBlocks,
219-
const int nThreads);
203+
o2::its::ExternalAllocator* alloc);
220204

221205
template <int NLayers = 7>
222206
void computeTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
@@ -237,9 +221,7 @@ void computeTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
237221
const bool shiftRefToCluster,
238222
const o2::base::Propagator* propagator,
239223
const o2::base::PropagatorF::MatCorrType matCorrType,
240-
o2::its::ExternalAllocator* alloc,
241-
const int nBlocks,
242-
const int nThreads);
224+
o2::its::ExternalAllocator* alloc);
243225

244226
} // namespace o2::its
245227
#endif // ITSTRACKINGGPU_TRACKINGKERNELS_H_

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -627,8 +627,7 @@ constexpr auto makeIterTags(std::index_sequence<I...>)
627627
{
628628
return std::array<uint64_t, sizeof...(I)>{makeIterTag<I>()...};
629629
}
630-
// FIXME: we have to be careful that the MaxIter does not diverge from the 4 here!
631-
constexpr auto kIterTags = makeIterTags(std::make_index_sequence<4>{});
630+
constexpr auto kIterTags = makeIterTags(std::make_index_sequence<constants::MaxIter>{});
632631
} // namespace detail
633632

634633
template <int NLayers>

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx‎

Lines changed: 3 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,6 @@
1818

1919
#include "ITStrackingGPU/TrackerTraitsGPU.h"
2020
#include "ITStrackingGPU/TrackingKernels.h"
21-
#include "ITStracking/TrackingConfigParam.h"
2221
#include "ITStracking/Constants.h"
2322

2423
namespace o2::its
@@ -63,8 +62,6 @@ void TrackerTraitsGPU<NLayers>::adoptTimeFrame(TimeFrame<NLayers>* tf)
6362
template <int NLayers>
6463
void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int iVertex)
6564
{
66-
const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
67-
6865
// start by queuing loading needed of two last layers
6966
for (int iLayer{NLayers}; iLayer-- > NLayers - 2;) {
7067
mTimeFrameGPU->createUsedClustersDevice(iteration, iLayer);
@@ -109,8 +106,6 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
109106
this->mTrkParams[iteration].LayerRadii,
110107
mTimeFrameGPU->getMSangles(),
111108
mTimeFrameGPU->getFrameworkAllocator(),
112-
conf.nBlocksLayerTracklets[iteration],
113-
conf.nThreadsLayerTracklets[iteration],
114109
mTimeFrameGPU->getStreams());
115110
mTimeFrameGPU->createTrackletsBuffers(iLayer);
116111
if (mTimeFrameGPU->getNTracklets()[iLayer] == 0) {
@@ -144,17 +139,13 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
144139
this->mTrkParams[iteration].LayerRadii,
145140
mTimeFrameGPU->getMSangles(),
146141
mTimeFrameGPU->getFrameworkAllocator(),
147-
conf.nBlocksLayerTracklets[iteration],
148-
conf.nThreadsLayerTracklets[iteration],
149142
mTimeFrameGPU->getStreams());
150143
}
151144
}
152145

153146
template <int NLayers>
154147
void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
155148
{
156-
auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
157-
158149
// start by queuing loading needed of three last layers
159150
for (int iLayer{NLayers}; iLayer-- > NLayers - 3;) {
160151
mTimeFrameGPU->loadUnsortedClustersDevice(iteration, iLayer);
@@ -194,8 +185,6 @@ void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
194185
this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
195186
this->mTrkParams[iteration].NSigmaCut,
196187
mTimeFrameGPU->getFrameworkAllocator(),
197-
conf.nBlocksLayerCells[iteration],
198-
conf.nThreadsLayerCells[iteration],
199188
mTimeFrameGPU->getStreams());
200189
mTimeFrameGPU->createCellsBuffers(iLayer);
201190
if (mTimeFrameGPU->getNCells()[iLayer] == 0) {
@@ -215,17 +204,13 @@ void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
215204
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
216205
this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
217206
this->mTrkParams[iteration].NSigmaCut,
218-
conf.nBlocksLayerCells[iteration],
219-
conf.nThreadsLayerCells[iteration],
220207
mTimeFrameGPU->getStreams());
221208
}
222209
}
223210

224211
template <int NLayers>
225212
void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
226213
{
227-
const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
228-
229214
for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NeighboursPerRoad(); ++iLayer) {
230215
const int currentLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer])};
231216
const int nextLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer + 1])};
@@ -248,8 +233,6 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
248233
nextLayerCellsNum,
249234
1e2,
250235
mTimeFrameGPU->getFrameworkAllocator(),
251-
conf.nBlocksFindNeighbours[iteration],
252-
conf.nThreadsFindNeighbours[iteration],
253236
mTimeFrameGPU->getStream(iLayer));
254237
mTimeFrameGPU->createNeighboursDevice(iLayer);
255238
if (mTimeFrameGPU->getNNeighbours()[iLayer] == 0) {
@@ -267,8 +250,6 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
267250
currentLayerCellsNum,
268251
nextLayerCellsNum,
269252
1e2,
270-
conf.nBlocksFindNeighbours[iteration],
271-
conf.nThreadsFindNeighbours[iteration],
272253
mTimeFrameGPU->getStream(iLayer));
273254
mTimeFrameGPU->getArrayNNeighbours()[iLayer] = filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
274255
mTimeFrameGPU->getDeviceNeighbours(iLayer),
@@ -282,7 +263,6 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
282263
template <int NLayers>
283264
void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
284265
{
285-
auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
286266
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
287267
const int minimumLayer{startLevel - 1};
288268
bounded_vector<CellSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
@@ -305,9 +285,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
305285
this->mTrkParams[0].MaxChi2NDF,
306286
mTimeFrameGPU->getDevicePropagator(),
307287
this->mTrkParams[0].CorrType,
308-
mTimeFrameGPU->getFrameworkAllocator(),
309-
conf.nBlocksProcessNeighbours[iteration],
310-
conf.nThreadsProcessNeighbours[iteration]);
288+
mTimeFrameGPU->getFrameworkAllocator());
311289
}
312290
// 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.
313291
if (trackSeeds.empty()) {
@@ -334,9 +312,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
334312
this->mTrkParams[0].ShiftRefToCluster,
335313
mTimeFrameGPU->getDevicePropagator(),
336314
this->mTrkParams[0].CorrType,
337-
mTimeFrameGPU->getFrameworkAllocator(),
338-
conf.nBlocksTracksSeeds[iteration],
339-
conf.nThreadsTracksSeeds[iteration]);
315+
mTimeFrameGPU->getFrameworkAllocator());
340316
mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size());
341317
computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
342318
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
@@ -356,9 +332,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
356332
this->mTrkParams[0].ShiftRefToCluster,
357333
mTimeFrameGPU->getDevicePropagator(),
358334
this->mTrkParams[0].CorrType,
359-
mTimeFrameGPU->getFrameworkAllocator(),
360-
conf.nBlocksTracksSeeds[iteration],
361-
conf.nThreadsTracksSeeds[iteration]);
335+
mTimeFrameGPU->getFrameworkAllocator());
362336
mTimeFrameGPU->downloadTrackITSExtDevice();
363337

364338
auto& tracks = mTimeFrameGPU->getTrackITSExt();

0 commit comments

Comments
 (0)