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