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 <algorithm>
16#include <array>
17
18#include "ITStrackingGPU/TrackerTraitsGPU.h"
19#include "ITStrackingGPU/TrackingKernels.h"
20#include "ITStracking/Configuration.h"
21
22namespace o2::its
23{
24template <int NLayers>
25void TrackerTraitsGPU<NLayers>::initialiseTimeFrame(const int iteration)
26{
27 mTimeFrameGPU->initialise(this->mTrkParams[iteration], NLayers, iteration);
28
29 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
30 // on default stream
31 mTimeFrameGPU->loadVertices();
32 // TODO these tables can be put in persistent memory
33 mTimeFrameGPU->loadROFOverlapTable(); // this can be put in constant memory actually
34 mTimeFrameGPU->loadROFVertexLookupTable();
35 mTimeFrameGPU->loadTrackingTopologies();
36 // once the tables are in persistent memory just update the vertex one
37 // mTimeFrameGPU->updateROFVertexLookupTable();
38 mTimeFrameGPU->loadIndexTableUtils();
39 // pinned on host
40 mTimeFrameGPU->createUsedClustersDeviceArray();
41 mTimeFrameGPU->createClustersDeviceArray();
42 mTimeFrameGPU->createUnsortedClustersDeviceArray();
43 mTimeFrameGPU->createClustersIndexTablesArray();
44 mTimeFrameGPU->createTrackingFrameInfoDeviceArray();
45 mTimeFrameGPU->createROFrameClustersDeviceArray();
46 // device array
47 mTimeFrameGPU->createTrackletsLUTDeviceArray();
48 mTimeFrameGPU->createTrackletsBuffersArray();
49 mTimeFrameGPU->createCellsBuffersArray();
50 mTimeFrameGPU->createCellsLUTDeviceArray();
51 }
52 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass] || this->mTrkParams[iteration].PassFlags[IterationStep::UseUPCMask]) {
53 mTimeFrameGPU->loadROFCutMask(iteration);
54 }
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 const bool loadFirstPassData = this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass] && iVertex <= 0; // load data only on first pass and first vertex
70 for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) {
71 if (loadFirstPassData) {
72 mTimeFrameGPU->createUsedClustersDevice(iLayer);
73 mTimeFrameGPU->loadClustersDevice(iLayer);
74 mTimeFrameGPU->loadClustersIndexTables(iLayer);
75 mTimeFrameGPU->loadROFrameClustersDevice(iLayer);
76 }
77 mTimeFrameGPU->recordEvent(iLayer);
78 }
79
80 for (int linkId{0}; linkId < hostTopology.nLinks; ++linkId) {
81 mTimeFrameGPU->createTrackletsLUTDevice(loadFirstPassData, linkId); // on first pass allocates, then only clears memory
82 }
83
84 // Stack allocations created from trackleting through road finding are scoped to one tracker pass.
85 // With per-primary-vertex processing, the chain is called once per vertex while initialisation is only done once.
86 mTimeFrameGPU->pushMemoryStack(iteration);
87
88 for (int linkId{0}; linkId < hostTopology.nLinks; ++linkId) {
89 const auto link = hostTopology.getLink(linkId);
90 mTimeFrameGPU->waitEvent(linkId, link.fromLayer);
91 mTimeFrameGPU->waitEvent(linkId, link.toLayer);
92 countTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
93 mTimeFrameGPU->getDeviceROFMaskTableView(),
94 linkId,
95 link.fromLayer,
96 link.toLayer,
97 mTimeFrameGPU->getDeviceROFOverlapTableView(),
98 mTimeFrameGPU->getDeviceROFVertexLookupTableView(),
99 iVertex,
100 mTimeFrameGPU->getDeviceVertices(),
101 mTimeFrameGPU->getDeviceROFramesPV(),
102 mTimeFrameGPU->getDeviceArrayClusters(),
103 mTimeFrameGPU->getClusterSizes(),
104 mTimeFrameGPU->getDeviceROFrameClusters(),
105 (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
106 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
107 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
108 mTimeFrameGPU->getDeviceTrackletsLUTs(),
109 this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices],
110 this->mTrkParams[iteration].NSigmaCut,
111 topology,
112 mTimeFrameGPU->getLinkPhiCuts(),
113 this->mTrkParams[iteration].PVres,
114 mTimeFrameGPU->getMinRs(),
115 mTimeFrameGPU->getMaxRs(),
116 mTimeFrameGPU->getPositionResolutions(),
117 this->mTrkParams[iteration].LayerRadii,
118 mTimeFrameGPU->getLinkMSAngles(),
119 mTimeFrameGPU->getFrameworkAllocator(),
120 mTimeFrameGPU->getStreams());
121 mTimeFrameGPU->createTrackletsBuffers(linkId);
122 if (mTimeFrameGPU->getNTracklets()[linkId] == 0) {
123 mTimeFrameGPU->recordEvent(linkId);
124 continue;
125 }
126 computeTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
127 mTimeFrameGPU->getDeviceROFMaskTableView(),
128 linkId,
129 link.fromLayer,
130 link.toLayer,
131 mTimeFrameGPU->getDeviceROFOverlapTableView(),
132 mTimeFrameGPU->getDeviceROFVertexLookupTableView(),
133 iVertex,
134 mTimeFrameGPU->getDeviceVertices(),
135 mTimeFrameGPU->getDeviceROFramesPV(),
136 mTimeFrameGPU->getDeviceArrayClusters(),
137 mTimeFrameGPU->getClusterSizes(),
138 mTimeFrameGPU->getDeviceROFrameClusters(),
139 (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
140 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
141 mTimeFrameGPU->getDeviceArrayTracklets(),
142 mTimeFrameGPU->getDeviceTracklets(),
143 mTimeFrameGPU->getNTracklets(),
144 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
145 mTimeFrameGPU->getDeviceTrackletsLUTs(),
146 this->mTrkParams[iteration].PassFlags[IterationStep::SelectUPCVertices],
147 this->mTrkParams[iteration].NSigmaCut,
148 topology,
149 mTimeFrameGPU->getLinkPhiCuts(),
150 this->mTrkParams[iteration].PVres,
151 mTimeFrameGPU->getMinRs(),
152 mTimeFrameGPU->getMaxRs(),
153 mTimeFrameGPU->getPositionResolutions(),
154 this->mTrkParams[iteration].LayerRadii,
155 mTimeFrameGPU->getLinkMSAngles(),
156 mTimeFrameGPU->getFrameworkAllocator(),
157 mTimeFrameGPU->getStreams());
158 mTimeFrameGPU->recordEvent(linkId);
159 }
160}
161
162template <int NLayers>
163void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
164{
165 const auto topology = mTimeFrameGPU->getDeviceTrackingTopologyView();
166 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
167 for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) {
168 if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
169 mTimeFrameGPU->loadUnsortedClustersDevice(iLayer);
170 mTimeFrameGPU->loadTrackingFrameInfoDevice(iLayer);
171 }
172 mTimeFrameGPU->recordEvent(iLayer);
173 }
174
175 for (int cellTopologyId{hostTopology.nCells}; cellTopologyId--;) {
176 const auto cellTopology = hostTopology.getCell(cellTopologyId);
177 const auto first = hostTopology.getLink(cellTopology.firstLink);
178 const auto second = hostTopology.getLink(cellTopology.secondLink);
179 const int currentLayerTrackletsNum{static_cast<int>(mTimeFrameGPU->getNTracklets()[cellTopology.firstLink])};
180 if (!currentLayerTrackletsNum || !mTimeFrameGPU->getNTracklets()[cellTopology.secondLink]) {
181 mTimeFrameGPU->getNCells()[cellTopologyId] = 0;
182 continue;
183 }
184
185 mTimeFrameGPU->createCellsLUTDevice(cellTopologyId);
186 mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.firstLink);
187 mTimeFrameGPU->waitEvent(cellTopologyId, cellTopology.secondLink);
188 mTimeFrameGPU->waitEvent(cellTopologyId, first.fromLayer);
189 mTimeFrameGPU->waitEvent(cellTopologyId, first.toLayer);
190 mTimeFrameGPU->waitEvent(cellTopologyId, second.toLayer);
191 countCellsHandler<NLayers>(mTimeFrameGPU->getDeviceArrayClusters(),
192 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
193 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
194 mTimeFrameGPU->getDeviceArrayTracklets(),
195 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
196 currentLayerTrackletsNum,
197 cellTopologyId,
198 topology,
199 nullptr,
200 mTimeFrameGPU->getDeviceArrayCellsLUT(),
201 mTimeFrameGPU->getDeviceCellLUTs()[cellTopologyId],
202 this->mBz,
203 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
204 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
205 this->mTrkParams[iteration].NSigmaCut,
206 this->mTrkParams[iteration].LayerxX0,
207 mTimeFrameGPU->getFrameworkAllocator(),
208 mTimeFrameGPU->getStreams());
209 mTimeFrameGPU->createCellsBuffers(cellTopologyId);
210 if (mTimeFrameGPU->getNCells()[cellTopologyId] == 0) {
211 mTimeFrameGPU->recordEvent(cellTopologyId);
212 continue;
213 }
214 computeCellsHandler<NLayers>(mTimeFrameGPU->getDeviceArrayClusters(),
215 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
216 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
217 mTimeFrameGPU->getDeviceArrayTracklets(),
218 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
219 currentLayerTrackletsNum,
220 cellTopologyId,
221 topology,
222 mTimeFrameGPU->getDeviceCells()[cellTopologyId],
223 mTimeFrameGPU->getDeviceArrayCellsLUT(),
224 mTimeFrameGPU->getDeviceCellLUTs()[cellTopologyId],
225 this->mBz,
226 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
227 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
228 this->mTrkParams[iteration].NSigmaCut,
229 this->mTrkParams[iteration].LayerxX0,
230 mTimeFrameGPU->getStreams());
231 mTimeFrameGPU->recordEvent(cellTopologyId);
232 }
233 mTimeFrameGPU->syncStreams(false);
234}
235
236template <int NLayers>
237void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
238{
239 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
240 for (int outerLayer{0}; outerLayer < NLayers; ++outerLayer) {
241 for (int targetCellTopologyId{0}; targetCellTopologyId < hostTopology.nCells; ++targetCellTopologyId) {
242 const auto targetCellTopology = hostTopology.getCell(targetCellTopologyId);
243 if (targetCellTopology.hitLayerMask.last() != outerLayer) {
244 continue;
245 }
246 const int targetCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[targetCellTopologyId])};
247 if (!targetCellsNum) {
248 mTimeFrameGPU->getNNeighbours()[targetCellTopologyId] = 0;
249 mTimeFrameGPU->recordEvent(targetCellTopologyId);
250 continue;
251 }
252 mTimeFrameGPU->createNeighboursIndexTablesDevice(targetCellTopologyId);
253 mTimeFrameGPU->createNeighboursLUTDevice(targetCellTopologyId, targetCellsNum);
254
255 for (int sourceCellTopologyId{0}; sourceCellTopologyId < hostTopology.nCells; ++sourceCellTopologyId) {
256 const auto sourceCellTopology = hostTopology.getCell(sourceCellTopologyId);
257 const int sourceCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[sourceCellTopologyId])};
258 if (!sourceCellsNum || sourceCellTopology.secondLink != targetCellTopology.firstLink) {
259 continue;
260 }
261 mTimeFrameGPU->waitEvent(targetCellTopologyId, sourceCellTopologyId);
262 countCellNeighboursHandler<NLayers>(mTimeFrameGPU->getDeviceArrayCells(),
263 mTimeFrameGPU->getDeviceNeighboursIndexTables(targetCellTopologyId),
264 mTimeFrameGPU->getDeviceArrayCellsLUT(),
265 sourceCellTopologyId,
266 targetCellTopologyId,
267 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
268 this->mBz,
269 sourceCellsNum,
270 mTimeFrameGPU->getStream(targetCellTopologyId));
271 }
272
273 scanCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighboursIndexTables(targetCellTopologyId),
274 mTimeFrameGPU->getDeviceNeighboursLUT(targetCellTopologyId),
275 targetCellsNum,
276 mTimeFrameGPU->getFrameworkAllocator(),
277 mTimeFrameGPU->getStream(targetCellTopologyId));
278
279 mTimeFrameGPU->createNeighboursDevice(targetCellTopologyId);
280 if (mTimeFrameGPU->getNNeighbours()[targetCellTopologyId] == 0) {
281 mTimeFrameGPU->recordEvent(targetCellTopologyId);
282 continue;
283 }
284
285 for (int sourceCellTopologyId{0}; sourceCellTopologyId < hostTopology.nCells; ++sourceCellTopologyId) {
286 const auto sourceCellTopology = hostTopology.getCell(sourceCellTopologyId);
287 const int sourceCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[sourceCellTopologyId])};
288 if (!sourceCellsNum || sourceCellTopology.secondLink != targetCellTopology.firstLink) {
289 continue;
290 }
291 computeCellNeighboursHandler<NLayers>(mTimeFrameGPU->getDeviceArrayCells(),
292 mTimeFrameGPU->getDeviceNeighboursIndexTables(targetCellTopologyId),
293 mTimeFrameGPU->getDeviceArrayCellsLUT(),
294 mTimeFrameGPU->getDeviceNeighbours(targetCellTopologyId),
295 sourceCellTopologyId,
296 targetCellTopologyId,
297 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
298 this->mBz,
299 sourceCellsNum,
300 mTimeFrameGPU->getStream(targetCellTopologyId));
301 }
302 mTimeFrameGPU->recordEvent(targetCellTopologyId);
303 }
304 }
305 mTimeFrameGPU->syncStreams(false);
306}
307
308template <int NLayers>
309void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
310{
311 bounded_vector<bounded_vector<int>> firstClusters(this->mTrkParams[iteration].NLayers, bounded_vector<int>(this->getMemoryPool().get()), this->getMemoryPool().get());
312 firstClusters.resize(this->mTrkParams[iteration].NLayers);
313 const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
314 const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop];
315 const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot];
316 const bool extendTracks = extendTop || extendBot;
317 for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
318 bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
319 for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) {
320 const int startLayer = hostTopology.getCell(startCellTopologyId).hitLayerMask.last();
321 if (!(this->mTrkParams[iteration].StartLayerMask.has(startLayer)) || mTimeFrameGPU->getNCells()[startCellTopologyId] == 0) {
322 continue;
323 }
324 processNeighboursHandler<NLayers>(startLevel,
325 startCellTopologyId,
326 mTimeFrameGPU->getDeviceArrayCells(),
327 mTimeFrameGPU->getDeviceCells()[startCellTopologyId],
328 nullptr,
329 nullptr,
330 mTimeFrameGPU->getArrayNCells().data(),
331 (const uint8_t**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
332 mTimeFrameGPU->getDeviceArrayNeighbours(),
333 mTimeFrameGPU->getDeviceArrayNeighboursCellLUT(),
334 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
335 trackSeeds,
336 this->mBz,
337 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
338 this->mTrkParams[iteration].MaxChi2NDF,
339 this->mTrkParams[iteration].MaxHoles,
340 this->mTrkParams[iteration].getMinSeedingClusters(),
341 this->mTrkParams[iteration].HoleLayerMask,
342 this->mTrkParams[iteration].getNonSeedingLayerMask(),
343 this->mTrkParams[iteration].LayerxX0,
344 mTimeFrameGPU->getDevicePropagator(),
345 this->mTrkParams[iteration].CorrType,
346 mTimeFrameGPU->getFrameworkAllocator());
347 }
348 // 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.
349 if (trackSeeds.empty()) {
350 LOGP(debug, "No track seeds found, skipping track finding");
351 continue;
352 }
353 mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);
354
355 // Since TrackITSExt is an enourmous class it is better to first count how many
356 // successfull fits we do and only then allocate
357 countTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
358 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
359 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
360 mTimeFrameGPU->getDeviceTrackSeedsLUT(),
361 this->mTrkParams[iteration].LayerRadii,
362 this->mTrkParams[iteration].MinPt,
363 this->mTrkParams[iteration].LayerxX0,
364 trackSeeds.size(),
365 this->mBz,
366 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
367 this->mTrkParams[iteration].MaxChi2NDF,
368 this->mTrkParams[iteration].ReseedIfShorter,
369 this->mTrkParams[iteration].RepeatRefitOut,
370 this->mTrkParams[iteration].ShiftRefToCluster,
371 mTimeFrameGPU->getDevicePropagator(),
372 this->mTrkParams[iteration].CorrType,
373 mTimeFrameGPU->getFrameworkAllocator());
374 mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size());
375 if (extendTracks) {
376 mTimeFrameGPU->createTrackExtensionScratchDevice(constants::GPUThreadsTotal, this->mTrkParams[iteration].TrackFollowerMaxHypotheses);
377 }
378 computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
379 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
380 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
381 mTimeFrameGPU->getDeviceIndexTableUtils(),
382 mTimeFrameGPU->getDeviceROFMaskTableView(),
383 mTimeFrameGPU->getDeviceROFOverlapTableView(),
384 mTimeFrameGPU->getDeviceArrayClusters(),
385 (const unsigned char**)mTimeFrameGPU->getDeviceArrayUsedClusters(),
386 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
387 mTimeFrameGPU->getDeviceROFrameClusters(),
388 mTimeFrameGPU->getDeviceTrackITSExt(),
389 mTimeFrameGPU->getDeviceTrackIndices(),
390 mTimeFrameGPU->getDeviceTrackSeedsLUT(),
391 extendTracks ? mTimeFrameGPU->getDeviceActiveTrackExtensionHypotheses() : nullptr,
392 extendTracks ? mTimeFrameGPU->getDeviceNextTrackExtensionHypotheses() : nullptr,
393 this->mTrkParams[iteration].LayerRadii,
394 this->mTrkParams[iteration].MinPt,
395 this->mTrkParams[iteration].LayerxX0,
396 trackSeeds.size(),
397 mTimeFrameGPU->getNTrackSeeds(),
398 this->mBz,
399 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
400 this->mTrkParams[iteration].MaxChi2NDF,
401 this->mTrkParams[iteration].ReseedIfShorter,
402 this->mTrkParams[iteration].RepeatRefitOut,
403 this->mTrkParams[iteration].ShiftRefToCluster,
404 this->mTrkParams[iteration].NLayers,
405 this->mTrkParams[iteration].PhiBins,
406 this->mTrkParams[iteration].TrackFollowerMaxHypotheses,
407 extendTop,
408 extendBot,
409 this->mTrkParams[iteration].TrackFollowerNSigmaCutPhi,
410 this->mTrkParams[iteration].TrackFollowerNSigmaCutZ,
411 mTimeFrameGPU->getDevicePropagator(),
412 this->mTrkParams[iteration].CorrType,
413 mTimeFrameGPU->getFrameworkAllocator());
414 mTimeFrameGPU->downloadTrackITSExtDevice();
415 mTimeFrameGPU->downloadTrackIndicesDevice();
416
417 auto& tracks = mTimeFrameGPU->getTrackITSExt();
418 const auto& trackIndices = mTimeFrameGPU->getTrackIndices();
419 this->acceptTracks(iteration, tracks, trackIndices, firstClusters);
420 mTimeFrameGPU->loadUsedClustersDevice();
421 }
422 this->markTracks(iteration);
423 // wipe the artefact memory
424 mTimeFrameGPU->popMemoryStack(iteration);
425};
426
427template <int NLayers>
428int TrackerTraitsGPU<NLayers>::getTFNumberOfClusters() const
429{
430 return mTimeFrameGPU->getNumberOfClusters();
431}
432
433template <int NLayers>
434int TrackerTraitsGPU<NLayers>::getTFNumberOfTracklets() const
435{
436 return std::accumulate(mTimeFrameGPU->getNTracklets().begin(), mTimeFrameGPU->getNTracklets().end(), 0);
437}
438
439template <int NLayers>
440int TrackerTraitsGPU<NLayers>::getTFNumberOfCells() const
441{
442 return mTimeFrameGPU->getNumberOfCells();
443}
444
445template <int NLayers>
446void TrackerTraitsGPU<NLayers>::setBz(float bz)
447{
448 this->mBz = bz;
449 mTimeFrameGPU->setBz(bz);
450}
451
452template class TrackerTraitsGPU<7>;
453#ifdef ENABLE_UPGRADES
454template class TrackerTraitsGPU<11>;
455template class TrackerTraitsGPU<13>;
456#endif
457} // 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:428
o2::its::TrackerTraitsGPU::initialiseTimeFrame
void initialiseTimeFrame(const int iteration) final
Definition TrackerTraitsGPU.cxx:25
o2::its::TrackerTraitsGPU::computeLayerCells
void computeLayerCells(const int iteration) final
Definition TrackerTraitsGPU.cxx:163
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:309
o2::its::TrackerTraitsGPU::getTFNumberOfCells
int getTFNumberOfCells() const override
Definition TrackerTraitsGPU.cxx:440
o2::its::TrackerTraitsGPU::findCellsNeighbours
void findCellsNeighbours(const int iteration) final
Definition TrackerTraitsGPU.cxx:237
o2::its::TrackerTraitsGPU::getTFNumberOfTracklets
int getTFNumberOfTracklets() const override
Definition TrackerTraitsGPU.cxx:434
o2::header::get
auto get(const std::byte *buffer, size_t=0)
Definition DataHeader.h:454
o2::its::constants::GPUThreadsTotal
constexpr int GPUThreadsTotal
Definition Constants.h:44
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 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:246
o2::its::scanCellNeighboursHandler
void scanCellNeighboursHandler(int *neighboursCursor, int *neighboursLUT, const unsigned int nCells, o2::its::ExternalAllocator *alloc, gpu::Stream &stream)
o2::its::extendBot
const bool const bool extendBot
Definition TrackFollower.h:122
o2::its::extendTop
const bool extendTop
Definition TrackFollower.h:121
o2::its::computeTrackSeedHandler
void computeTrackSeedHandler(TrackSeed< NLayers > *trackSeeds, const TrackingFrameInfo **foundTrackingFrameInfo, const Cluster **unsortedClusters, const IndexTableUtils< NLayers > *utils, const typename ROFMaskTable< NLayers >::View &rofMask, const typename ROFOverlapTable< NLayers >::View &rofOverlaps, const Cluster **clusters, const unsigned char **usedClusters, const int **clustersIndexTables, const int **ROFClusters, o2::its::TrackITSExt *tracks, int *trackIndices, const int *seedLUT, TrackExtensionHypothesis< NLayers > *activeHypotheses, TrackExtensionHypothesis< NLayers > *nextHypotheses, 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 float maxChi2ClusterAttachment, const float maxChi2NDF, const int reseedIfShorter, const bool repeatRefitOut, const bool shiftRefToCluster, const int nLayers, const int phiBins, const int maxHypotheses, const bool extendTop, const bool extendBot, const float nSigmaCutPhi, const float nSigmaCutZ, 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