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#include "ITStracking/Constants.h"
22
23namespace o2::its
24{
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 const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
52
53 int startROF{this->mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * this->mTrkParams[iteration].nROFsPerIterations : 0};
54 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())};
55
56 mTimeFrameGPU->createTrackletsLUTDevice(iteration);
57 countTrackletsInROFsHandler<nLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
58 mTimeFrameGPU->getDeviceMultCutMask(),
59 startROF,
60 endROF,
61 mTimeFrameGPU->getNrof(),
62 this->mTrkParams[iteration].DeltaROF,
63 iVertex,
64 mTimeFrameGPU->getDeviceVertices(),
65 mTimeFrameGPU->getDeviceROFramesPV(),
66 mTimeFrameGPU->getPrimaryVerticesNum(),
67 mTimeFrameGPU->getDeviceArrayClusters(),
68 mTimeFrameGPU->getClusterSizes(),
69 mTimeFrameGPU->getDeviceROframeClusters(),
70 mTimeFrameGPU->getDeviceArrayUsedClusters(),
71 mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
72 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
73 mTimeFrameGPU->getDeviceTrackletsLUTs(), // Required for the exclusive sums
74 iteration,
75 this->mTrkParams[iteration].NSigmaCut,
76 mTimeFrameGPU->getPhiCuts(),
77 this->mTrkParams[iteration].PVres,
78 mTimeFrameGPU->getMinRs(),
79 mTimeFrameGPU->getMaxRs(),
80 mTimeFrameGPU->getPositionResolutions(),
81 this->mTrkParams[iteration].LayerRadii,
82 mTimeFrameGPU->getMSangles(),
83 conf.nBlocksLayerTracklets[iteration],
84 conf.nThreadsLayerTracklets[iteration],
85 mTimeFrameGPU->getStreams());
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.nBlocksLayerTracklets[iteration],
117 conf.nThreadsLayerTracklets[iteration],
118 mTimeFrameGPU->getStreams());
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 mTimeFrameGPU->getNCells()[iLayer] = 0;
130 continue;
131 }
132 const int currentLayerTrackletsNum{static_cast<int>(mTimeFrameGPU->getNTracklets()[iLayer])};
133 countCellsHandler(mTimeFrameGPU->getDeviceArrayClusters(),
134 mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
135 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
136 mTimeFrameGPU->getDeviceArrayTracklets(),
137 mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
138 mTimeFrameGPU->getNTracklets()[iLayer],
139 iLayer,
140 nullptr,
141 mTimeFrameGPU->getDeviceArrayCellsLUT(),
142 mTimeFrameGPU->getDeviceCellLUTs()[iLayer],
143 this->mTrkParams[iteration].DeltaROF,
144 this->mBz,
145 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
146 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
147 this->mTrkParams[iteration].NSigmaCut,
148 conf.nBlocksLayerCells[iteration],
149 conf.nThreadsLayerCells[iteration]);
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->mTrkParams[iteration].DeltaROF,
162 this->mBz,
163 this->mTrkParams[iteration].MaxChi2ClusterAttachment,
164 this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
165 this->mTrkParams[iteration].NSigmaCut,
166 conf.nBlocksLayerCells[iteration],
167 conf.nThreadsLayerCells[iteration]);
168 }
169}
170
171template <int nLayers>
172void TrackerTraitsGPU<nLayers>::findCellsNeighbours(const int iteration)
173{
174 mTimeFrameGPU->createNeighboursIndexTablesDevice();
175 const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
176 for (int iLayer{0}; iLayer < this->mTrkParams[iteration].CellsPerRoad() - 1; ++iLayer) {
177 const int currentLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer])};
178 const int nextLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer + 1])};
179 if (!nextLayerCellsNum || !currentLayerCellsNum) {
180 mTimeFrameGPU->getNNeighbours()[iLayer] = 0;
181 continue;
182 }
183
184 mTimeFrameGPU->createNeighboursLUTDevice(iLayer, nextLayerCellsNum);
185 unsigned int nNeigh = countCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
186 mTimeFrameGPU->getDeviceNeighboursLUT(iLayer), // LUT is initialised here.
187 mTimeFrameGPU->getDeviceArrayCellsLUT(),
188 mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
189 mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer),
190 (const Tracklet**)mTimeFrameGPU->getDeviceArrayTracklets(),
191 this->mTrkParams[0].DeltaROF,
192 this->mTrkParams[0].MaxChi2ClusterAttachment,
193 this->mBz,
194 iLayer,
195 currentLayerCellsNum,
196 nextLayerCellsNum,
197 1e2,
198 conf.nBlocksFindNeighbours[iteration],
199 conf.nThreadsFindNeighbours[iteration]);
200
201 mTimeFrameGPU->createNeighboursDevice(iLayer, nNeigh);
202
203 computeCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
204 mTimeFrameGPU->getDeviceNeighboursLUT(iLayer),
205 mTimeFrameGPU->getDeviceArrayCellsLUT(),
206 mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
207 mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer),
208 (const Tracklet**)mTimeFrameGPU->getDeviceArrayTracklets(),
209 this->mTrkParams[0].DeltaROF,
210 this->mTrkParams[0].MaxChi2ClusterAttachment,
211 this->mBz,
212 iLayer,
213 currentLayerCellsNum,
214 nextLayerCellsNum,
215 1e2,
216 conf.nBlocksFindNeighbours[iteration],
217 conf.nThreadsFindNeighbours[iteration]);
218
219 nNeigh = filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
220 mTimeFrameGPU->getDeviceNeighbours(iLayer),
221 nNeigh,
222 mTimeFrameGPU->getExternalAllocator());
223 mTimeFrameGPU->getArrayNNeighbours()[iLayer] = nNeigh;
224 }
225 mTimeFrameGPU->createNeighboursDeviceArray();
226 mTimeFrameGPU->unregisterRest();
227};
228
229template <int nLayers>
230void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
231{
232 auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
233 for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
234 const int minimumLayer{startLevel - 1};
235 bounded_vector<CellSeed> trackSeeds(this->getMemoryPool().get());
236 for (int startLayer{this->mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) {
237 if ((this->mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) {
238 continue;
239 }
240 processNeighboursHandler<nLayers>(startLayer,
241 startLevel,
242 mTimeFrameGPU->getDeviceArrayCells(),
243 mTimeFrameGPU->getDeviceCells()[startLayer],
244 mTimeFrameGPU->getArrayNCells(),
245 mTimeFrameGPU->getDeviceArrayUsedClusters(),
246 mTimeFrameGPU->getDeviceNeighboursAll(),
247 mTimeFrameGPU->getDeviceNeighboursLUTs(),
248 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
249 trackSeeds,
250 mTimeFrameGPU->getExternalAllocator(),
251 this->mBz,
252 this->mTrkParams[0].MaxChi2ClusterAttachment,
253 this->mTrkParams[0].MaxChi2NDF,
254 mTimeFrameGPU->getDevicePropagator(),
255 this->mTrkParams[0].CorrType,
256 conf.nBlocksProcessNeighbours[iteration],
257 conf.nThreadsProcessNeighbours[iteration]);
258 }
259 // 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.
260 if (trackSeeds.empty()) {
261 LOGP(debug, "No track seeds found, skipping track finding");
262 continue;
263 }
264 mTimeFrameGPU->createTrackITSExtDevice(trackSeeds);
265 mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);
266
267 trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds
268 mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo
269 mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks
270 this->mTrkParams[iteration].MinPt, // std::vector<float>& minPtsHost,
271 trackSeeds.size(), // const size_t nSeeds
272 this->mBz, // const float Bz
273 startLevel, // const int startLevel,
274 this->mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment
275 this->mTrkParams[0].MaxChi2NDF, // float maxChi2NDF
276 mTimeFrameGPU->getDevicePropagator(), // const o2::base::Propagator* propagator
277 this->mTrkParams[0].CorrType, // o2::base::PropagatorImpl<float>::MatCorrType
278 conf.nBlocksTracksSeeds[iteration],
279 conf.nThreadsTracksSeeds[iteration]);
280
281 mTimeFrameGPU->downloadTrackITSExtDevice(trackSeeds);
282
283 auto& tracks = mTimeFrameGPU->getTrackITSExt();
284
285 for (auto& track : tracks) {
286 if (!track.getChi2()) {
287 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.
288 }
289 int nShared = 0;
290 bool isFirstShared{false};
291 for (int iLayer{0}; iLayer < this->mTrkParams[0].NLayers; ++iLayer) {
292 if (track.getClusterIndex(iLayer) == constants::UnusedIndex) {
293 continue;
294 }
295 nShared += int(mTimeFrameGPU->isClusterUsed(iLayer, track.getClusterIndex(iLayer)));
296 isFirstShared |= !iLayer && mTimeFrameGPU->isClusterUsed(iLayer, track.getClusterIndex(iLayer));
297 }
298
299 if (nShared > this->mTrkParams[0].ClusterSharing) {
300 continue;
301 }
302
303 std::array<int, 3> rofs{INT_MAX, INT_MAX, INT_MAX};
304 for (int iLayer{0}; iLayer < this->mTrkParams[0].NLayers; ++iLayer) {
305 if (track.getClusterIndex(iLayer) == constants::UnusedIndex) {
306 continue;
307 }
308 mTimeFrameGPU->markUsedCluster(iLayer, track.getClusterIndex(iLayer));
309 int currentROF = mTimeFrameGPU->getClusterROF(iLayer, track.getClusterIndex(iLayer));
310 for (int iR{0}; iR < 3; ++iR) {
311 if (rofs[iR] == INT_MAX) {
312 rofs[iR] = currentROF;
313 }
314 if (rofs[iR] == currentROF) {
315 break;
316 }
317 }
318 }
319 if (rofs[2] != INT_MAX) {
320 continue;
321 }
322 if (rofs[1] != INT_MAX) {
323 track.setNextROFbit();
324 }
325 mTimeFrameGPU->getTracks(std::min(rofs[0], rofs[1])).emplace_back(track);
326 }
327 mTimeFrameGPU->loadUsedClustersDevice();
328 }
329};
330
331template <int nLayers>
332int TrackerTraitsGPU<nLayers>::getTFNumberOfClusters() const
333{
334 return mTimeFrameGPU->getNumberOfClusters();
335}
336
337template <int nLayers>
338int TrackerTraitsGPU<nLayers>::getTFNumberOfTracklets() const
339{
340 return std::accumulate(mTimeFrameGPU->getNTracklets().begin(), mTimeFrameGPU->getNTracklets().end(), 0);
341}
342
343template <int nLayers>
344int TrackerTraitsGPU<nLayers>::getTFNumberOfCells() const
345{
346 return mTimeFrameGPU->getNumberOfCells();
347}
348
349template <int nLayers>
350void TrackerTraitsGPU<nLayers>::setBz(float bz)
351{
352 this->mBz = bz;
353 mTimeFrameGPU->setBz(bz);
354}
355
356template class TrackerTraitsGPU<7>;
357} // namespace o2::its
TrackITS.h
Definition of the ITS track.
debug
std::ostringstream debug
Definition VariantJSONHelpers.cxx:305
o2::its::TrackerTraitsGPU::getTFNumberOfClusters
int getTFNumberOfClusters() const override
Definition TrackerTraitsGPU.cxx:332
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:172
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:338
o2::its::TrackerTraitsGPU::getTFNumberOfCells
int getTFNumberOfCells() const override
Definition TrackerTraitsGPU.cxx:344
o2::its::TrackerTraitsGPU::findRoads
void findRoads(const int iteration) final
Definition TrackerTraitsGPU.cxx:230
o2::header::get
auto get(const std::byte *buffer, size_t=0)
Definition DataHeader.h:454
o2::its::constants::UnusedIndex
constexpr int UnusedIndex
Definition Constants.h:30
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::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 int deltaROF, 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 Tracklet **tracklets, const int deltaROF, 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::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 int deltaROF, const float bz, const float maxChi2ClusterAttachment, const float cellDeltaTanLambdaSigma, const float nSigmaCut, 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 Tracklet **tracklets, const int deltaROF, 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)
tf
std::unique_ptr< GPUReconstructionTimeframe > tf
Definition standalone.cxx:84