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