Project
Loading...
Searching...
No Matches
TrackerTraitsGPU.cxx
Go to the documentation of this file.
1// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3// All rights not expressly granted are reserved.
4//
5// This software is distributed under the terms of the GNU General Public
6// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7//
8// In applying this license CERN does not waive the privileges and immunities
9// granted to it by virtue of its status as an Intergovernmental Organization
10// or submit itself to any jurisdiction.
12
13#include <unistd.h>
14
15#include "ITStrackingGPU/TrackerTraitsGPU.h"
16#include "ITStrackingGPU/TrackingKernels.h"
17#include "ITStracking/Configuration.h"
18
19namespace o2::its
20{
21
22template <int NLayers>
23void TrackerTraitsGPU<NLayers>::initialiseTimeFrame(const int iteration)
24{
25 mTimeFrameGPU->initialise(this->mTrkParams[iteration], NLayers, iteration);
26
27 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
28 // on default stream
29 mTimeFrameGPU->loadVertices();
30 // TODO these tables can be put in persistent memory
31 mTimeFrameGPU->loadROFOverlapTable(); // this can be put in constant memory actually
32 mTimeFrameGPU->loadROFVertexLookupTable();
33 mTimeFrameGPU->loadTrackingTopologies();
34 // once the tables are in persistent memory just update the vertex one
35 // mTimeFrameGPU->updateROFVertexLookupTable();
36 mTimeFrameGPU->loadIndexTableUtils();
37 // pinned on host
38 mTimeFrameGPU->createUsedClustersDeviceArray();
39 mTimeFrameGPU->createClustersDeviceArray();
40 mTimeFrameGPU->createUnsortedClustersDeviceArray();
41 mTimeFrameGPU->createClustersIndexTablesArray();
42 mTimeFrameGPU->createTrackingFrameInfoDeviceArray();
43 mTimeFrameGPU->createROFrameClustersDeviceArray();
44 // device array
45 mTimeFrameGPU->createTrackletsLUTDeviceArray();
46 mTimeFrameGPU->createTrackletsBuffersArray();
47 mTimeFrameGPU->createCellsBuffersArray();
48 mTimeFrameGPU->createCellsLUTDeviceArray();
49 }
50 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass] || this->mTrkParams[iteration].PassFlags[IterationStep::UseUPCMask]) {
51 mTimeFrameGPU->loadROFCutMask(iteration);
52 }
53 // push every create artefact on the stack
54 mTimeFrameGPU->pushMemoryStack(iteration);
55}
56
57template <int NLayers>
58void TrackerTraitsGPU<NLayers>::adoptTimeFrame(TimeFrame<NLayers>* tf)
59{
60 mTimeFrameGPU = static_cast<gpu::TimeFrameGPU<NLayers>*>(tf);
61 this->mTimeFrame = static_cast<TimeFrame<NLayers>*>(tf);
62}
63
64template <int NLayers>
65void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int iVertex)
66{
67 const auto topology = mTimeFrameGPU->getDeviceTrackingTopologyView();
68 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
69 for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) {
70 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
71 mTimeFrameGPU->createUsedClustersDevice(iLayer);
72 mTimeFrameGPU->loadClustersDevice(iLayer);
73 mTimeFrameGPU->loadClustersIndexTables(iLayer);
74 mTimeFrameGPU->loadROFrameClustersDevice(iLayer);
75 }
76 mTimeFrameGPU->recordEvent(iLayer);
77 }
78
79 for (int transitionId{0}; transitionId < hostTopology.nTransitions; ++transitionId) {
80 const auto transition = hostTopology.getTransition(transitionId);
81 mTimeFrameGPU->createTrackletsLUTDevice(this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass], transitionId);
82 mTimeFrameGPU->waitEvent(transitionId, transition.fromLayer);
83 mTimeFrameGPU->waitEvent(transitionId, transition.toLayer);
84 countTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
85 mTimeFrameGPU->getDeviceROFMaskTableView(),
86 transitionId,
87 transition.fromLayer,
88 transition.toLayer,
89 mTimeFrameGPU->getDeviceROFOverlapTableView(),
90 mTimeFrameGPU->getDeviceROFVertexLookupTableView(),
91 iVertex,
92 mTimeFrameGPU->getDeviceVertices(),
93 mTimeFrameGPU->getDeviceROFramesPV(),
94 mTimeFrameGPU->getDeviceArrayClusters(),
95 mTimeFrameGPU->getClusterSizes(),
96 mTimeFrameGPU->getDeviceROFrameClusters(),
97 (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
98 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
99 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
100 mTimeFrameGPU->getDeviceTrackletsLUTs(),
101 this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices],
102 this->mTrkParams[iteration].NSigmaCut,
103 topology,
104 mTimeFrameGPU->getTransitionPhiCuts(),
105 this->mTrkParams[iteration].PVres,
106 mTimeFrameGPU->getMinRs(),
107 mTimeFrameGPU->getMaxRs(),
108 mTimeFrameGPU->getPositionResolutions(),
109 this->mTrkParams[iteration].LayerRadii,
110 mTimeFrameGPU->getTransitionMSAngles(),
111 mTimeFrameGPU->getFrameworkAllocator(),
112 mTimeFrameGPU->getStreams());
113 mTimeFrameGPU->createTrackletsBuffers(transitionId);
114 if (mTimeFrameGPU->getNTracklets()[transitionId] == 0) {
115 mTimeFrameGPU->recordEvent(transitionId);
116 continue;
117 }
118 computeTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
119 mTimeFrameGPU->getDeviceROFMaskTableView(),
120 transitionId,
121 transition.fromLayer,
122 transition.toLayer,
123 mTimeFrameGPU->getDeviceROFOverlapTableView(),
124 mTimeFrameGPU->getDeviceROFVertexLookupTableView(),
125 iVertex,
126 mTimeFrameGPU->getDeviceVertices(),
127 mTimeFrameGPU->getDeviceROFramesPV(),
128 mTimeFrameGPU->getDeviceArrayClusters(),
129 mTimeFrameGPU->getClusterSizes(),
130 mTimeFrameGPU->getDeviceROFrameClusters(),
131 (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
132 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
133 mTimeFrameGPU->getDeviceArrayTracklets(),
134 mTimeFrameGPU->getDeviceTracklets(),
135 mTimeFrameGPU->getNTracklets(),
136 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
137 mTimeFrameGPU->getDeviceTrackletsLUTs(),
138 this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices],
139 this->mTrkParams[iteration].NSigmaCut,
140 topology,
141 mTimeFrameGPU->getTransitionPhiCuts(),
142 this->mTrkParams[iteration].PVres,
143 mTimeFrameGPU->getMinRs(),
144 mTimeFrameGPU->getMaxRs(),
145 mTimeFrameGPU->getPositionResolutions(),
146 this->mTrkParams[iteration].LayerRadii,
147 mTimeFrameGPU->getTransitionMSAngles(),
148 mTimeFrameGPU->getFrameworkAllocator(),
149 mTimeFrameGPU->getStreams());
150 mTimeFrameGPU->recordEvent(transitionId);
151 }
152}
153
154template <int NLayers>
155void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
156{
157 const auto topology = mTimeFrameGPU->getDeviceTrackingTopologyView();
158 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
159 for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) {
160 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
161 mTimeFrameGPU->loadUnsortedClustersDevice(iLayer);
162 mTimeFrameGPU->loadTrackingFrameInfoDevice(iLayer);
163 }
164 mTimeFrameGPU->recordEvent(iLayer);
165 }
166
167 for (int cellTopologyId{hostTopology.nCells}; cellTopologyId--;) {
168 const auto cellTopology = hostTopology.getCell(cellTopologyId);
169 const auto first = hostTopology.getTransition(cellTopology.firstTransition);
170 const auto second = hostTopology.getTransition(cellTopology.secondTransition);
171 const int currentLayerTrackletsNum{static_cast<int>(mTimeFrameGPU->getNTracklets()[cellTopology.firstTransition])};
172 if (!currentLayerTrackletsNum || !mTimeFrameGPU->getNTracklets()[cellTopology.secondTransition]) {
173 mTimeFrameGPU->getNCells()[cellTopologyId] = 0;
174 continue;
175 }
176
177 mTimeFrameGPU->createCellsLUTDevice(cellTopologyId);
178 mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.firstTransition);
179 mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.secondTransition);
180 mTimeFrameGPU->waitEvent(cellTopologyId, first.fromLayer);
181 mTimeFrameGPU->waitEvent(cellTopologyId, first.toLayer);
182 mTimeFrameGPU->waitEvent(cellTopologyId, second.toLayer);
183 countCellsHandler<NLayers>(mTimeFrameGPU->getDeviceArrayClusters(),
184 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
185 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
186 mTimeFrameGPU->getDeviceArrayTracklets(),
187 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
188 currentLayerTrackletsNum,
189 cellTopologyId,
190 topology,
191 nullptr,
192 mTimeFrameGPU->getDeviceArrayCellsLUT(),
193 mTimeFrameGPU->getDeviceCellLUTs()[cellTopologyId],
194 this->mBz,
195 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
196 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
197 this->mTrkParams[iteration].NSigmaCut,
198 this->mTrkParams[iteration].LayerxX0,
199 mTimeFrameGPU->getFrameworkAllocator(),
200 mTimeFrameGPU->getStreams());
201 mTimeFrameGPU->createCellsBuffers(cellTopologyId);
202 if (mTimeFrameGPU->getNCells()[cellTopologyId] == 0) {
203 mTimeFrameGPU->recordEvent(cellTopologyId);
204 continue;
205 }
206 computeCellsHandler<NLayers>(mTimeFrameGPU->getDeviceArrayClusters(),
207 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
208 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
209 mTimeFrameGPU->getDeviceArrayTracklets(),
210 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
211 currentLayerTrackletsNum,
212 cellTopologyId,
213 topology,
214 mTimeFrameGPU->getDeviceCells()[cellTopologyId],
215 mTimeFrameGPU->getDeviceArrayCellsLUT(),
216 mTimeFrameGPU->getDeviceCellLUTs()[cellTopologyId],
217 this->mBz,
218 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
219 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
220 this->mTrkParams[iteration].NSigmaCut,
221 this->mTrkParams[iteration].LayerxX0,
222 mTimeFrameGPU->getStreams());
223 mTimeFrameGPU->recordEvent(cellTopologyId);
224 }
225 mTimeFrameGPU->syncStreams(false);
226}
227
228template <int NLayers>
229void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
230{
231 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
232 for (int outerLayer{0}; outerLayer < NLayers; ++outerLayer) {
233 for (int targetCellTopologyId{0}; targetCellTopologyId < hostTopology.nCells; ++targetCellTopologyId) {
234 const auto targetCellTopology = hostTopology.getCell(targetCellTopologyId);
235 if (targetCellTopology.hitLayerMask.last() != outerLayer) {
236 continue;
237 }
238 const int targetCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[targetCellTopologyId])};
239 if (!targetCellsNum) {
240 mTimeFrameGPU->getNNeighbours()[targetCellTopologyId] = 0;
241 mTimeFrameGPU->recordEvent(targetCellTopologyId);
242 continue;
243 }
244 mTimeFrameGPU->createNeighboursIndexTablesDevice(targetCellTopologyId);
245 mTimeFrameGPU->createNeighboursLUTDevice(targetCellTopologyId, targetCellsNum);
246
247 for (int sourceCellTopologyId{0}; sourceCellTopologyId < hostTopology.nCells; ++sourceCellTopologyId) {
248 const auto sourceCellTopology = hostTopology.getCell(sourceCellTopologyId);
249 const int sourceCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[sourceCellTopologyId])};
250 if (!sourceCellsNum || sourceCellTopology.secondTransition != targetCellTopology.firstTransition) {
251 continue;
252 }
253 mTimeFrameGPU->waitEvent(targetCellTopologyId, sourceCellTopologyId);
254 countCellNeighboursHandler<NLayers>(mTimeFrameGPU->getDeviceArrayCells(),
255 mTimeFrameGPU->getDeviceNeighboursIndexTables(targetCellTopologyId),
256 mTimeFrameGPU->getDeviceArrayCellsLUT(),
257 sourceCellTopologyId,
258 targetCellTopologyId,
259 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
260 this->mBz,
261 sourceCellsNum,
262 mTimeFrameGPU->getStream(targetCellTopologyId));
263 }
264
265 scanCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighboursIndexTables(targetCellTopologyId),
266 mTimeFrameGPU->getDeviceNeighboursLUT(targetCellTopologyId),
267 targetCellsNum,
268 mTimeFrameGPU->getFrameworkAllocator(),
269 mTimeFrameGPU->getStream(targetCellTopologyId));
270
271 mTimeFrameGPU->createNeighboursDevice(targetCellTopologyId);
272 if (mTimeFrameGPU->getNNeighbours()[targetCellTopologyId] == 0) {
273 mTimeFrameGPU->recordEvent(targetCellTopologyId);
274 continue;
275 }
276
277 for (int sourceCellTopologyId{0}; sourceCellTopologyId < hostTopology.nCells; ++sourceCellTopologyId) {
278 const auto sourceCellTopology = hostTopology.getCell(sourceCellTopologyId);
279 const int sourceCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[sourceCellTopologyId])};
280 if (!sourceCellsNum || sourceCellTopology.secondTransition != targetCellTopology.firstTransition) {
281 continue;
282 }
283 computeCellNeighboursHandler<NLayers>(mTimeFrameGPU->getDeviceArrayCells(),
284 mTimeFrameGPU->getDeviceNeighboursIndexTables(targetCellTopologyId),
285 mTimeFrameGPU->getDeviceArrayCellsLUT(),
286 mTimeFrameGPU->getDeviceNeighbours(targetCellTopologyId),
287 sourceCellTopologyId,
288 targetCellTopologyId,
289 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
290 this->mBz,
291 sourceCellsNum,
292 mTimeFrameGPU->getStream(targetCellTopologyId));
293 }
294 mTimeFrameGPU->recordEvent(targetCellTopologyId);
295 }
296 }
297 mTimeFrameGPU->syncStreams(false);
298}
299
300template <int NLayers>
301void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
302{
303 bounded_vector<bounded_vector<int>> firstClusters(this->mTrkParams[iteration].NLayers, bounded_vector<int>(this->getMemoryPool().get()), this->getMemoryPool().get());
304 bounded_vector<bounded_vector<int>> sharedFirstClusters(this->mTrkParams[iteration].NLayers, bounded_vector<int>(this->getMemoryPool().get()), this->getMemoryPool().get());
305 firstClusters.resize(this->mTrkParams[iteration].NLayers);
306 sharedFirstClusters.resize(this->mTrkParams[iteration].NLayers);
307 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
308 for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
309 bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
310 for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) {
311 const int startLayer = hostTopology.getCell(startCellTopologyId).hitLayerMask.last();
312 if ((this->mTrkParams[iteration].StartLayerMask & (1 << startLayer)) == 0 ||
313 mTimeFrameGPU->getNCells()[startCellTopologyId] == 0) {
314 continue;
315 }
316 processNeighboursHandler<NLayers>(startLevel,
317 startCellTopologyId,
318 mTimeFrameGPU->getDeviceArrayCells(),
319 mTimeFrameGPU->getDeviceCells()[startCellTopologyId],
320 nullptr,
321 nullptr,
322 mTimeFrameGPU->getArrayNCells().data(),
323 (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
324 mTimeFrameGPU->getDeviceArrayNeighbours(),
325 mTimeFrameGPU->getDeviceArrayNeighboursCellLUT(),
326 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
327 trackSeeds,
328 this->mBz,
329 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
330 this->mTrkParams[iteration].MaxChi2NDF,
331 this->mTrkParams[iteration].MaxHoles,
332 this->mTrkParams[iteration].MinTrackLength,
333 this->mTrkParams[iteration].HoleLayerMask,
334 this->mTrkParams[iteration].LayerxX0,
335 mTimeFrameGPU->getDevicePropagator(),
336 this->mTrkParams[iteration].CorrType,
337 mTimeFrameGPU->getFrameworkAllocator());
338 }
339 // 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.
340 if (trackSeeds.empty()) {
341 LOGP(debug, "No track seeds found, skipping track finding");
342 continue;
343 }
344 mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);
345
346 // Since TrackITSExt is an enourmous class it is better to first count how many
347 // successfull fits we do and only then allocate
348 countTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
349 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
350 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
351 mTimeFrameGPU->getDeviceTrackSeedsLUT(),
352 this->mTrkParams[iteration].LayerRadii,
353 this->mTrkParams[iteration].MinPt,
354 this->mTrkParams[iteration].LayerxX0,
355 trackSeeds.size(),
356 this->mBz,
357 startLevel,
358 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
359 this->mTrkParams[iteration].MaxChi2NDF,
360 this->mTrkParams[iteration].ReseedIfShorter,
361 this->mTrkParams[iteration].RepeatRefitOut,
362 this->mTrkParams[iteration].ShiftRefToCluster,
363 mTimeFrameGPU->getDevicePropagator(),
364 this->mTrkParams[iteration].CorrType,
365 mTimeFrameGPU->getFrameworkAllocator());
366 mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size());
367 computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
368 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
369 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
370 mTimeFrameGPU->getDeviceTrackITSExt(),
371 mTimeFrameGPU->getDeviceTrackSeedsLUT(),
372 this->mTrkParams[iteration].LayerRadii,
373 this->mTrkParams[iteration].MinPt,
374 this->mTrkParams[iteration].LayerxX0,
375 trackSeeds.size(),
376 mTimeFrameGPU->getNTrackSeeds(),
377 this->mBz,
378 startLevel,
379 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
380 this->mTrkParams[iteration].MaxChi2NDF,
381 this->mTrkParams[iteration].ReseedIfShorter,
382 this->mTrkParams[iteration].RepeatRefitOut,
383 this->mTrkParams[iteration].ShiftRefToCluster,
384 mTimeFrameGPU->getDevicePropagator(),
385 this->mTrkParams[iteration].CorrType,
386 mTimeFrameGPU->getFrameworkAllocator());
387 mTimeFrameGPU->downloadTrackITSExtDevice();
388
389 auto& tracks = mTimeFrameGPU->getTrackITSExt();
390 this->acceptTracks(iteration, tracks, firstClusters, sharedFirstClusters);
391 mTimeFrameGPU->loadUsedClustersDevice();
392 }
393 this->markTracks(iteration, sharedFirstClusters);
394 // wipe the artefact memory
395 mTimeFrameGPU->popMemoryStack(iteration);
396};
397
398template <int NLayers>
399int TrackerTraitsGPU<NLayers>::getTFNumberOfClusters() const
400{
401 return mTimeFrameGPU->getNumberOfClusters();
402}
403
404template <int NLayers>
405int TrackerTraitsGPU<NLayers>::getTFNumberOfTracklets() const
406{
407 return std::accumulate(mTimeFrameGPU->getNTracklets().begin(), mTimeFrameGPU->getNTracklets().end(), 0);
408}
409
410template <int NLayers>
411int TrackerTraitsGPU<NLayers>::getTFNumberOfCells() const
412{
413 return mTimeFrameGPU->getNumberOfCells();
414}
415
416template <int NLayers>
417void TrackerTraitsGPU<NLayers>::setBz(float bz)
418{
419 this->mBz = bz;
420 mTimeFrameGPU->setBz(bz);
421}
422
423template class TrackerTraitsGPU<7>;
424#ifdef ENABLE_UPGRADES
425template class TrackerTraitsGPU<11>;
426#endif
427} // namespace o2::its
debug
std::ostringstream debug
Definition ExpressionJSONHelpers.cxx:83
o2::its::TrackerTraitsGPU::computeLayerTracklets
void computeLayerTracklets(const int iteration, int) final
Definition TrackerTraitsGPU.cxx:65
o2::its::TrackerTraitsGPU::getTFNumberOfClusters
int getTFNumberOfClusters() const override
Definition TrackerTraitsGPU.cxx:399
o2::its::TrackerTraitsGPU::initialiseTimeFrame
void initialiseTimeFrame(const int iteration) final
Definition TrackerTraitsGPU.cxx:23
o2::its::TrackerTraitsGPU::computeLayerCells
void computeLayerCells(const int iteration) final
Definition TrackerTraitsGPU.cxx:155
o2::its::TrackerTraitsGPU::adoptTimeFrame
void adoptTimeFrame(TimeFrame< NLayers > *tf) final
Definition TrackerTraitsGPU.cxx:58
o2::its::TrackerTraitsGPU::findRoads
void findRoads(const int iteration) final
Definition TrackerTraitsGPU.cxx:301
o2::its::TrackerTraitsGPU::getTFNumberOfCells
int getTFNumberOfCells() const override
Definition TrackerTraitsGPU.cxx:411
o2::its::TrackerTraitsGPU::findCellsNeighbours
void findCellsNeighbours(const int iteration) final
Definition TrackerTraitsGPU.cxx:229
o2::its::TrackerTraitsGPU::getTFNumberOfTracklets
int getTFNumberOfTracklets() const override
Definition TrackerTraitsGPU.cxx:405
o2::header::get
auto get(const std::byte *buffer, size_t=0)
Definition DataHeader.h:454
o2::its::countTrackSeedHandler
void countTrackSeedHandler(TrackSeed< NLayers > *trackSeeds, const TrackingFrameInfo **foundTrackingFrameInfo, const Cluster **unsortedClusters, int *seedLUT, const std::vector< float > &layerRadiiHost, const std::vector< float > &minPtsHost, const std::vector< float > &layerxX0Host, const unsigned int nSeeds, const float Bz, const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, const o2::base::Propagator *propagator, const o2::base::PropagatorF::MatCorrType matCorrType, o2::its::ExternalAllocator *alloc)
o2::its::bounded_vector
std::pmr::vector< T > bounded_vector
Definition BoundedAllocator.h:219
o2::its::scanCellNeighboursHandler
void scanCellNeighboursHandler(int *neighboursCursor, int *neighboursLUT, const unsigned int nCells, o2::its::ExternalAllocator *alloc, gpu::Stream &stream)
o2::its::computeTrackSeedHandler
void computeTrackSeedHandler(TrackSeed< NLayers > *trackSeeds, const TrackingFrameInfo **foundTrackingFrameInfo, const Cluster **unsortedClusters, o2::its::TrackITSExt *tracks, const int *seedLUT, const std::vector< float > &layerRadiiHost, const std::vector< float > &minPtsHost, const std::vector< float > &layerxX0Host, const unsigned int nSeeds, const unsigned int nTracks, const float Bz, const int startLevel, const float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, const o2::base::Propagator *propagator, const o2::base::PropagatorF::MatCorrType matCorrType, o2::its::ExternalAllocator *alloc)
tf
std::unique_ptr< GPUReconstructionTimeframe > tf
Definition standalone.cxx:86
o2::its::TimeFrame::getTrackingTopologyView
const auto & getTrackingTopologyView() const
Definition TimeFrame.h:141