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