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