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 <array>
14#include <unistd.h>
15
16#include "DataFormatsITS/TrackITS.h"
17
18#include "ITStrackingGPU/TrackerTraitsGPU.h"
19#include "ITStrackingGPU/TrackingKernels.h"
20#include "ITStracking/TrackingConfigParam.h"
21namespace o2::its
22{
23constexpr int UnusedIndex{-1};
24
25template <int nLayers>
26void TrackerTraitsGPU<nLayers>::initialiseTimeFrame(const int iteration)
27{
28 mTimeFrameGPU->initialise(iteration, this->mTrkParams[iteration], nLayers);
29 mTimeFrameGPU->loadClustersDevice(iteration);
30 mTimeFrameGPU->loadUnsortedClustersDevice(iteration);
31 mTimeFrameGPU->loadClustersIndexTables(iteration);
32 mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration);
33 mTimeFrameGPU->loadMultiplicityCutMask(iteration);
34 mTimeFrameGPU->loadVertices(iteration);
35 mTimeFrameGPU->loadROframeClustersDevice(iteration);
36 mTimeFrameGPU->createUsedClustersDevice(iteration);
37 mTimeFrameGPU->loadIndexTableUtils(iteration);
38}
39
40template <int nLayers>
41void TrackerTraitsGPU<nLayers>::adoptTimeFrame(TimeFrame<nLayers>* tf)
42{
43 mTimeFrameGPU = static_cast<gpu::TimeFrameGPU<nLayers>*>(tf);
44 this->mTimeFrame = static_cast<TimeFrame<nLayers>*>(tf);
45}
46
47template <int nLayers>
48void TrackerTraitsGPU<nLayers>::computeLayerTracklets(const int iteration, int iROFslice, int iVertex)
49{
50 auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
51 mTimeFrameGPU->createTrackletsLUTDevice(iteration);
52
53 const Vertex diamondVert({this->mTrkParams[iteration].Diamond[0], this->mTrkParams[iteration].Diamond[1], this->mTrkParams[iteration].Diamond[2]}, {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f}, 1, 1.f);
54 gsl::span<const Vertex> diamondSpan(&diamondVert, 1);
55 int startROF{this->mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * this->mTrkParams[iteration].nROFsPerIterations : 0};
56 int endROF{o2::gpu::CAMath::Min(this->mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * this->mTrkParams[iteration].nROFsPerIterations + this->mTrkParams[iteration].DeltaROF : mTimeFrameGPU->getNrof(), mTimeFrameGPU->getNrof())};
57
58 countTrackletsInROFsHandler<nLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
59 mTimeFrameGPU->getDeviceMultCutMask(),
60 startROF,
61 endROF,
62 mTimeFrameGPU->getNrof(),
63 this->mTrkParams[iteration].DeltaROF,
64 iVertex,
65 mTimeFrameGPU->getDeviceVertices(),
66 mTimeFrameGPU->getDeviceROFramesPV(),
67 mTimeFrameGPU->getPrimaryVerticesNum(),
68 mTimeFrameGPU->getDeviceArrayClusters(),
69 mTimeFrameGPU->getClusterSizes(),
70 mTimeFrameGPU->getDeviceROframeClusters(),
71 mTimeFrameGPU->getDeviceArrayUsedClusters(),
72 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
73 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
74 mTimeFrameGPU->getDeviceTrackletsLUTs(), // Required for the exclusive sums
75 iteration,
76 this->mTrkParams[iteration].NSigmaCut,
77 mTimeFrameGPU->getPhiCuts(),
78 this->mTrkParams[iteration].PVres,
79 mTimeFrameGPU->getMinRs(),
80 mTimeFrameGPU->getMaxRs(),
81 mTimeFrameGPU->getPositionResolutions(),
82 this->mTrkParams[iteration].LayerRadii,
83 mTimeFrameGPU->getMSangles(),
84 conf.nBlocks,
85 conf.nThreads);
86 mTimeFrameGPU->createTrackletsBuffers();
87 computeTrackletsInROFsHandler<nLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
88 mTimeFrameGPU->getDeviceMultCutMask(),
89 startROF,
90 endROF,
91 mTimeFrameGPU->getNrof(),
92 this->mTrkParams[iteration].DeltaROF,
93 iVertex,
94 mTimeFrameGPU->getDeviceVertices(),
95 mTimeFrameGPU->getDeviceROFramesPV(),
96 mTimeFrameGPU->getPrimaryVerticesNum(),
97 mTimeFrameGPU->getDeviceArrayClusters(),
98 mTimeFrameGPU->getClusterSizes(),
99 mTimeFrameGPU->getDeviceROframeClusters(),
100 mTimeFrameGPU->getDeviceArrayUsedClusters(),
101 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
102 mTimeFrameGPU->getDeviceArrayTracklets(),
103 mTimeFrameGPU->getDeviceTracklet(),
104 mTimeFrameGPU->getNTracklets(),
105 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
106 mTimeFrameGPU->getDeviceTrackletsLUTs(),
107 iteration,
108 this->mTrkParams[iteration].NSigmaCut,
109 mTimeFrameGPU->getPhiCuts(),
110 this->mTrkParams[iteration].PVres,
111 mTimeFrameGPU->getMinRs(),
112 mTimeFrameGPU->getMaxRs(),
113 mTimeFrameGPU->getPositionResolutions(),
114 this->mTrkParams[iteration].LayerRadii,
115 mTimeFrameGPU->getMSangles(),
116 conf.nBlocks,
117 conf.nThreads);
118}
119
120template <int nLayers>
121void TrackerTraitsGPU<nLayers>::computeLayerCells(const int iteration)
122{
123 mTimeFrameGPU->createCellsLUTDevice();
124 auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
125
126 for (int iLayer = 0; iLayer < this->mTrkParams[iteration].CellsPerRoad(); ++iLayer) {
127 if (!mTimeFrameGPU->getNTracklets()[iLayer + 1] || !mTimeFrameGPU->getNTracklets()[iLayer]) {
128 continue;
129 }
130 const int currentLayerTrackletsNum{static_cast<int>(mTimeFrameGPU->getNTracklets()[iLayer])};
131 countCellsHandler(mTimeFrameGPU->getDeviceArrayClusters(),
132 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
133 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
134 mTimeFrameGPU->getDeviceArrayTracklets(),
135 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
136 mTimeFrameGPU->getNTracklets()[iLayer],
137 iLayer,
138 nullptr,
139 mTimeFrameGPU->getDeviceArrayCellsLUT(),
140 mTimeFrameGPU->getDeviceCellLUTs()[iLayer],
141 this->mBz,
142 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
143 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
144 this->mTrkParams[iteration].NSigmaCut,
145 conf.nBlocks,
146 conf.nThreads);
147 mTimeFrameGPU->createCellsBuffers(iLayer);
148 computeCellsHandler(mTimeFrameGPU->getDeviceArrayClusters(),
149 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
150 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
151 mTimeFrameGPU->getDeviceArrayTracklets(),
152 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
153 mTimeFrameGPU->getNTracklets()[iLayer],
154 iLayer,
155 mTimeFrameGPU->getDeviceCells()[iLayer],
156 mTimeFrameGPU->getDeviceArrayCellsLUT(),
157 mTimeFrameGPU->getDeviceCellLUTs()[iLayer],
158 this->mBz,
159 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
160 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
161 this->mTrkParams[iteration].NSigmaCut,
162 conf.nBlocks,
163 conf.nThreads);
164 }
165}
166
167template <int nLayers>
168void TrackerTraitsGPU<nLayers>::findCellsNeighbours(const int iteration)
169{
170 mTimeFrameGPU->createNeighboursIndexTablesDevice();
171 auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
172 for (int iLayer{0}; iLayer < this->mTrkParams[iteration].CellsPerRoad() - 1; ++iLayer) {
173 const int nextLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer + 1])};
174
175 if (!nextLayerCellsNum) {
176 continue;
177 }
178
179 mTimeFrameGPU->createNeighboursLUTDevice(iLayer, nextLayerCellsNum);
180 unsigned int nNeigh = countCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
181 mTimeFrameGPU->getDeviceNeighboursLUT(iLayer), // LUT is initialised here.
182 mTimeFrameGPU->getDeviceArrayCellsLUT(),
183 mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
184 mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer),
185 this->mTrkParams[0].MaxChi2ClusterAttachment,
186 this->mBz,
187 iLayer,
188 mTimeFrameGPU->getNCells()[iLayer],
189 nextLayerCellsNum,
190 1e2,
191 conf.nBlocks,
192 conf.nThreads);
193
194 mTimeFrameGPU->createNeighboursDevice(iLayer, nNeigh);
195
196 computeCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
197 mTimeFrameGPU->getDeviceNeighboursLUT(iLayer),
198 mTimeFrameGPU->getDeviceArrayCellsLUT(),
199 mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
200 mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer),
201 this->mTrkParams[0].MaxChi2ClusterAttachment,
202 this->mBz,
203 iLayer,
204 mTimeFrameGPU->getNCells()[iLayer],
205 nextLayerCellsNum,
206 1e2,
207 conf.nBlocks,
208 conf.nThreads);
209
210 filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
211 mTimeFrameGPU->getDeviceNeighbours(iLayer),
212 nNeigh);
213 }
214 mTimeFrameGPU->createNeighboursDeviceArray();
215 mTimeFrameGPU->unregisterRest();
216};
217
218template <int nLayers>
219void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
220{
221 auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
222 for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
223 const int minimumLayer{startLevel - 1};
224 bounded_vector<CellSeed> trackSeeds(this->getMemoryPool().get());
225 for (int startLayer{this->mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) {
226 if ((this->mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) {
227 continue;
228 }
229 processNeighboursHandler<nLayers>(startLayer,
230 startLevel,
231 mTimeFrameGPU->getDeviceArrayCells(),
232 mTimeFrameGPU->getDeviceCells()[startLayer],
233 mTimeFrameGPU->getArrayNCells(),
234 mTimeFrameGPU->getDeviceArrayUsedClusters(),
235 mTimeFrameGPU->getDeviceNeighboursAll(),
236 mTimeFrameGPU->getDeviceNeighboursLUTs(),
237 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
238 trackSeeds,
239 this->mBz,
240 this->mTrkParams[0].MaxChi2ClusterAttachment,
241 this->mTrkParams[0].MaxChi2NDF,
242 mTimeFrameGPU->getDevicePropagator(),
243 this->mCorrType,
244 conf.nBlocks,
245 conf.nThreads);
246 }
247 // 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.
248 if (!trackSeeds.size()) {
249 LOGP(info, "No track seeds found, skipping track finding");
250 continue;
251 }
252 mTimeFrameGPU->createTrackITSExtDevice(trackSeeds);
253 mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);
254
255 trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds
256 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo
257 mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks
258 this->mTrkParams[iteration].MinPt, // std::vector<float>& minPtsHost,
259 trackSeeds.size(), // const size_t nSeeds
260 this->mBz, // const float Bz
261 startLevel, // const int startLevel,
262 this->mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment
263 this->mTrkParams[0].MaxChi2NDF, // float maxChi2NDF
264 mTimeFrameGPU->getDevicePropagator(), // const o2::base::Propagator* propagator
265 this->mCorrType, // o2::base::PropagatorImpl<float>::MatCorrType
266 conf.nBlocks,
267 conf.nThreads);
268
269 mTimeFrameGPU->downloadTrackITSExtDevice(trackSeeds);
270
271 auto& tracks = mTimeFrameGPU->getTrackITSExt();
272
273 for (auto& track : tracks) {
274 if (!track.getChi2()) {
275 continue; // this is to skip the unset tracks that are put at the beginning of the vector by the sorting. To see if this can be optimised.
276 }
277 int nShared = 0;
278 bool isFirstShared{false};
279 for (int iLayer{0}; iLayer < this->mTrkParams[0].NLayers; ++iLayer) {
280 if (track.getClusterIndex(iLayer) == UnusedIndex) {
281 continue;
282 }
283 nShared += int(mTimeFrameGPU->isClusterUsed(iLayer, track.getClusterIndex(iLayer)));
284 isFirstShared |= !iLayer && mTimeFrameGPU->isClusterUsed(iLayer, track.getClusterIndex(iLayer));
285 }
286
287 if (nShared > this->mTrkParams[0].ClusterSharing) {
288 continue;
289 }
290
291 std::array<int, 3> rofs{INT_MAX, INT_MAX, INT_MAX};
292 for (int iLayer{0}; iLayer < this->mTrkParams[0].NLayers; ++iLayer) {
293 if (track.getClusterIndex(iLayer) == UnusedIndex) {
294 continue;
295 }
296 mTimeFrameGPU->markUsedCluster(iLayer, track.getClusterIndex(iLayer));
297 int currentROF = mTimeFrameGPU->getClusterROF(iLayer, track.getClusterIndex(iLayer));
298 for (int iR{0}; iR < 3; ++iR) {
299 if (rofs[iR] == INT_MAX) {
300 rofs[iR] = currentROF;
301 }
302 if (rofs[iR] == currentROF) {
303 break;
304 }
305 }
306 }
307 if (rofs[2] != INT_MAX) {
308 continue;
309 }
310 if (rofs[1] != INT_MAX) {
311 track.setNextROFbit();
312 }
313 mTimeFrameGPU->getTracks(std::min(rofs[0], rofs[1])).emplace_back(track);
314 }
315 mTimeFrameGPU->loadUsedClustersDevice();
316 }
317 if (iteration == this->mTrkParams.size() - 1) {
318 mTimeFrameGPU->unregisterHostMemory(0);
319 }
320};
321
322template <int nLayers>
323int TrackerTraitsGPU<nLayers>::getTFNumberOfClusters() const
324{
325 return mTimeFrameGPU->getNumberOfClusters();
326}
327
328template <int nLayers>
329int TrackerTraitsGPU<nLayers>::getTFNumberOfTracklets() const
330{
331 return std::accumulate(mTimeFrameGPU->getNTracklets().begin(), mTimeFrameGPU->getNTracklets().end(), 0);
332}
333
334template <int nLayers>
335int TrackerTraitsGPU<nLayers>::getTFNumberOfCells() const
336{
337 return mTimeFrameGPU->getNumberOfCells();
338}
339
340template <int nLayers>
341void TrackerTraitsGPU<nLayers>::setBz(float bz)
342{
343 this->mBz = bz;
344 mTimeFrameGPU->setBz(bz);
345}
346
347template class TrackerTraitsGPU<7>;
348} // namespace o2::its
TrackITS.h
Definition of the ITS track.
o2::its::TrackerTraitsGPU::getTFNumberOfClusters
int getTFNumberOfClusters() const override
Definition TrackerTraitsGPU.cxx:323
o2::its::TrackerTraitsGPU::initialiseTimeFrame
void initialiseTimeFrame(const int iteration) final
Definition TrackerTraitsGPU.cxx:26
o2::its::TrackerTraitsGPU::findCellsNeighbours
void findCellsNeighbours(const int iteration) final
Definition TrackerTraitsGPU.cxx:168
o2::its::TrackerTraitsGPU::adoptTimeFrame
void adoptTimeFrame(TimeFrame< nLayers > *tf) final
Definition TrackerTraitsGPU.cxx:41
o2::its::TrackerTraitsGPU::computeLayerCells
void computeLayerCells(const int iteration) final
Definition TrackerTraitsGPU.cxx:121
o2::its::TrackerTraitsGPU::computeLayerTracklets
void computeLayerTracklets(const int iteration, int, int) final
Definition TrackerTraitsGPU.cxx:48
o2::its::TrackerTraitsGPU::getTFNumberOfTracklets
int getTFNumberOfTracklets() const override
Definition TrackerTraitsGPU.cxx:329
o2::its::TrackerTraitsGPU::getTFNumberOfCells
int getTFNumberOfCells() const override
Definition TrackerTraitsGPU.cxx:335
o2::its::TrackerTraitsGPU::findRoads
void findRoads(const int iteration) final
Definition TrackerTraitsGPU.cxx:219
o2::header::get
auto get(const std::byte *buffer, size_t=0)
Definition DataHeader.h:454
o2::its::filterCellNeighboursHandler
int filterCellNeighboursHandler(gpuPair< int, int > *, int *, unsigned int)
o2::its::bounded_vector
std::pmr::vector< T > bounded_vector
Definition BoundedAllocator.h:116
o2::its::UnusedIndex
constexpr int UnusedIndex
Definition TrackerTraitsGPU.cxx:23
o2::its::computeCellsHandler
void computeCellsHandler(const Cluster **sortedClusters, const Cluster **unsortedClusters, const TrackingFrameInfo **tfInfo, Tracklet **tracklets, int **trackletsLUT, const int nTracklets, const int layer, CellSeed *cells, int **cellsLUTsDeviceArray, int *cellsLUTsHost, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, const float nSigmaCut, const int nBlocks, const int nThreads)
o2::its::computeCellNeighboursHandler
void computeCellNeighboursHandler(CellSeed **cellsLayersDevice, int *neighboursLUTs, int **cellsLUTs, gpuPair< int, int > *cellNeighbours, int *neighboursIndexTable, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, const unsigned int nCells, const unsigned int nCellsNext, const int maxCellNeighbours, const int nBlocks, const int nThreads)
o2::its::countCellNeighboursHandler
unsigned int countCellNeighboursHandler(CellSeed **cellsLayersDevice, int *neighboursLUTs, int **cellsLUTs, gpuPair< int, int > *cellNeighbours, int *neighboursIndexTable, const float maxChi2ClusterAttachment, const float bz, const int layerIndex, const unsigned int nCells, const unsigned int nCellsNext, const int maxCellNeighbours, const int nBlocks, const int nThreads)
o2::its::trackSeedHandler
void trackSeedHandler(CellSeed *trackSeeds, const TrackingFrameInfo **foundTrackingFrameInfo, o2::its::TrackITSExt *tracks, std::vector< float > &minPtsHost, const unsigned int nSeeds, const float Bz, const int startLevel, float maxChi2ClusterAttachment, float maxChi2NDF, const o2::base::Propagator *propagator, const o2::base::PropagatorF::MatCorrType matCorrType, const int nBlocks, const int nThreads)
o2::its::countCellsHandler
void countCellsHandler(const Cluster **sortedClusters, const Cluster **unsortedClusters, const TrackingFrameInfo **tfInfo, Tracklet **tracklets, int **trackletsLUT, const int nTracklets, const int layer, CellSeed *cells, int **cellsLUTsDeviceArray, int *cellsLUTsHost, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, const float nSigmaCut, const int nBlocks, const int nThreads)
tf
std::unique_ptr< GPUReconstructionTimeframe > tf
Definition standalone.cxx:84