Project
Loading...
Searching...
No Matches
GPUTPCTrackletConstructor.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.
11
14
15#define GPUCA_CADEBUG 0
16
17#include "GPUTPCDef.h"
18#include "GPUTPCGrid.h"
19#include "GPUTPCHit.h"
20#include "GPUTPCTrackParam.h"
21#include "GPUTPCTracker.h"
22#include "GPUTPCTracklet.h"
26#include "CalibdEdxContainer.h"
27#include "GPUParam.inc"
28#include "GPUCommonMath.h"
29
30using namespace o2::gpu;
31
33{
34 // Initialize Tracklet Parameters using default values
35 tParam.InitParam();
36}
37
39{
40 bool ok = 1;
41 const float* c = tParam.Cov();
42 for (int32_t i = 0; i < 15; i++) {
43 ok = ok && CAMath::Finite(c[i]);
44 }
45 for (int32_t i = 0; i < 5; i++) {
46 ok = ok && CAMath::Finite(tParam.Par()[i]);
47 }
48 ok = ok && (tParam.X() > 50);
49 if (c[0] <= 0 || c[2] <= 0 || c[5] <= 0 || c[9] <= 0 || c[14] <= 0) {
50 ok = 0;
51 }
52 return (ok);
53}
54
55GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUTPCTrackParam& GPUrestrict() tParam, calink* rowHits)
56{
57 // reconstruction of tracklets, tracklet store step
58 const uint32_t nHits = r.mLastRow + 1 - r.mFirstRow;
59 if (nHits == 0 || r.mNHits == 0 || (r.mNHits < GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(tParam.QPt() * tracker.Param().qptB5Scaler) || !CheckCov(tParam) || CAMath::Abs(tParam.GetQPt() * tracker.Param().qptB5Scaler) > tracker.Param().rec.maxTrackQPtB5)) {
60 CADEBUG(printf(" Rejected: nHits %d QPt %f MinHits %d MaxQPt %f CheckCov %d\n", r.mNHits, tParam.QPt(), GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(tParam.QPt() * tracker.Param().qptB5Scaler), tracker.Param().rec.maxTrackQPtB5, (int32_t)CheckCov(tParam)));
61 return;
62 }
63
64 /*GPUInfo("Tracklet %d: Hits %3d NDF %3d Chi %8.4f Sign %f Cov: %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f %2.4f", r.mISH, r.mNHits, tParam.GetNDF(), tParam.GetChi2(), tParam.GetSignCosPhi(),
65 tParam.Cov()[0], tParam.Cov()[1], tParam.Cov()[2], tParam.Cov()[3], tParam.Cov()[4], tParam.Cov()[5], tParam.Cov()[6], tParam.Cov()[7], tParam.Cov()[8], tParam.Cov()[9],
66 tParam.Cov()[10], tParam.Cov()[11], tParam.Cov()[12], tParam.Cov()[13], tParam.Cov()[14]);*/
67
68 uint32_t hitout = CAMath::AtomicAdd(tracker.NRowHits(), nHits);
69 if (hitout + nHits > tracker.NMaxRowHits()) {
70 tracker.raiseError(GPUErrors::ERROR_TRACKLET_HIT_OVERFLOW, tracker.ISector(), hitout + nHits, tracker.NMaxRowHits());
71 CAMath::AtomicExch(tracker.NRowHits(), tracker.NMaxRowHits());
72 return;
73 }
74 uint32_t itrout = CAMath::AtomicAdd(tracker.NTracklets(), 1u);
75 if (itrout >= tracker.NMaxTracklets()) {
76 tracker.raiseError(GPUErrors::ERROR_TRACKLET_OVERFLOW, tracker.ISector(), itrout, tracker.NMaxTracklets());
77 CAMath::AtomicExch(tracker.NTracklets(), tracker.NMaxTracklets());
78 return;
79 }
80
81 GPUglobalref() GPUTPCTracklet& GPUrestrict() tracklet = tracker.Tracklets()[itrout];
82
83 CADEBUG(printf(" Storing tracklet: %d rows\n", nHits));
84
85 tracklet.SetFirstRow(r.mFirstRow);
86 tracklet.SetLastRow(r.mLastRow);
87 tracklet.SetFirstHit(hitout);
88 tracklet.SetParam(tParam.GetParam());
89 int32_t w = tracker.CalculateHitWeight(r.mNHits, tParam.GetChi2());
90 tracklet.SetHitWeight(w);
91#ifdef __HIPCC__ // Todo: fixme!
92 for (int32_t iRow = r.mFirstRow - 1; ++iRow <= r.mLastRow; /*iRow++*/) {
93#else
94 for (int32_t iRow = r.mFirstRow; iRow <= r.mLastRow; iRow++) {
95#endif
96 calink ih = rowHits[iRow];
97 tracker.TrackletRowHits()[hitout + (iRow - r.mFirstRow)] = ih;
98 if (ih != CALINK_INVAL && ih != CALINK_DEAD_CHANNEL) {
99 CA_MAKE_SHARED_REF(GPUTPCRow, row, tracker.Row(iRow), s.mRows[iRow]);
100 tracker.MaximizeHitWeight(row, ih, w);
101 }
102 }
103}
104
105template <class T>
106GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, int32_t /*nThreads*/, int32_t /*iBlock*/, int32_t /*iThread*/, GPUsharedref() T& s, GPUTPCThreadMemory& GPUrestrict() r, GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUTPCTrackParam& GPUrestrict() tParam, int32_t iRow, calink& rowHit, calink* rowHits)
107{
108 // reconstruction of tracklets, tracklets update step
109 CA_MAKE_SHARED_REF(GPUTPCRow, row, tracker.Row(iRow), s.mRows[iRow]);
110
111 float y0 = row.Grid().YMin();
112 float stepY = row.HstepY();
113 float z0 = row.Grid().ZMin() - tParam.ZOffset();
114 float stepZ = row.HstepZ();
115
116 if (r.mStage == 0) { // fitting part
117 do {
118 if (iRow < r.mStartRow || r.mCurrIH == CALINK_INVAL) {
119 break;
120 }
121 if ((iRow - r.mStartRow) & 1) {
122 rowHit = CALINK_INVAL;
123 break; // SG!!! - jump over the row
124 }
125
126 cahit2 hh = tracker.HitData(row)[r.mCurrIH];
127
128 int32_t seedIH = r.mCurrIH;
129 r.mCurrIH = tracker.HitLinkUpData(row)[r.mCurrIH];
130
131 float x = row.X();
132 float y = y0 + hh.x * stepY;
133 float z = z0 + hh.y * stepZ;
134 if (iRow != r.mStartRow || !tracker.Param().par.continuousTracking) {
135 tParam.ConstrainZ(z, tracker.ISector(), z0, r.mLastZ);
136 tracker.GetConstantMem()->calibObjects.fastTransformHelper->TransformXYZ(tracker.ISector(), iRow, x, y, z);
137 }
138 if (iRow == r.mStartRow) {
139 if (tracker.Param().par.continuousTracking) {
140 float refZ = ((z > 0) ? tracker.Param().rec.tpc.defaultZOffsetOverR : -tracker.Param().rec.tpc.defaultZOffsetOverR) * x;
141 float zTmp = refZ;
142 tracker.GetConstantMem()->calibObjects.fastTransformHelper->TransformXYZ(tracker.ISector(), iRow, x, y, zTmp);
143 z += zTmp - refZ; // Add zCorrection (=zTmp - refZ) to z, such that zOffset is set such, that transformed (z - zOffset) becomes refZ
144 tParam.SetZOffset(z - refZ);
145 tParam.SetZ(refZ);
146 r.mLastZ = refZ;
147 } else {
148 tParam.SetZ(z);
149 r.mLastZ = z;
150 tParam.SetZOffset(0.f);
151 }
152 tParam.SetX(x);
153 tParam.SetY(y);
154 r.mLastY = y;
155 CADEBUG(printf("Tracklet %5d: FIT INIT ROW %3d X %8.3f -", r.mISH, iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
156 } else {
157 float dx = x - tParam.X();
158 float dy, dz;
159 if (r.mNHits >= 10) {
160 dy = y - tParam.Y();
161 dz = z - tParam.Z();
162 } else {
163 dy = y - r.mLastY;
164 dz = z - r.mLastZ;
165 }
166 r.mLastY = y;
167 r.mLastZ = z;
168
169 float ri = 1.f / CAMath::Sqrt(dx * dx + dy * dy);
170 if (iRow == r.mStartRow + 2) {
171 tParam.SetSinPhi(dy * ri);
172 tParam.SetSignCosPhi(dx);
173 tParam.SetDzDs(dz * ri);
174 float err2Y, err2Z;
175 tracker.GetErrors2Seeding(iRow, tParam, -1.f, err2Y, err2Z); // Use correct time
176 tParam.SetCov(0, err2Y);
177 tParam.SetCov(2, err2Z);
178 }
179 float sinPhi, cosPhi;
180 if (r.mNHits >= 10 && CAMath::Abs(tParam.SinPhi()) < GPUCA_MAX_SIN_PHI_LOW) {
181 sinPhi = tParam.SinPhi();
182 cosPhi = CAMath::Sqrt(1 - sinPhi * sinPhi);
183 } else {
184 sinPhi = dy * ri;
185 cosPhi = dx * ri;
186 }
187 CADEBUG(printf("%14s: FIT TRACK ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
188 if (!tParam.TransportToX(x, sinPhi, cosPhi, tracker.Param().bzCLight, GPUCA_MAX_SIN_PHI)) {
189 rowHit = CALINK_INVAL;
190 break;
191 }
192 CADEBUG(printf("%5s hits %3d: FIT PROP ROW %3d X %8.3f -", "", r.mNHits, iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
193 float err2Y, err2Z;
194 tracker.GetErrors2Seeding(iRow, tParam.GetZ(), sinPhi, tParam.GetDzDs(), -1.f, err2Y, err2Z); // TODO: Use correct time
195
196 if (r.mNHits >= 10) {
197 const float sErr2 = tracker.Param().GetSystematicClusterErrorIFC2(x, tParam.GetY(), tParam.GetZ(), tracker.ISector() >= 18);
198 err2Y += sErr2;
199 err2Z += sErr2;
200 const float kFactor = tracker.Param().rec.tpc.hitPickUpFactor * tracker.Param().rec.tpc.hitPickUpFactor * 3.5f * 3.5f;
201 float sy2 = kFactor * (tParam.Err2Y() + err2Y);
202 float sz2 = kFactor * (tParam.Err2Z() + err2Z);
203 if (sy2 > tracker.Param().rec.tpc.hitSearchArea2) {
204 sy2 = tracker.Param().rec.tpc.hitSearchArea2;
205 }
206 if (sz2 > tracker.Param().rec.tpc.hitSearchArea2) {
207 sz2 = tracker.Param().rec.tpc.hitSearchArea2;
208 }
209 dy = y - tParam.Y();
210 dz = z - tParam.Z();
211 if (dy * dy > sy2 || dz * dz > sz2) {
212 if (++r.mNMissed >= tracker.Param().rec.tpc.trackFollowingMaxRowGapSeed) {
213 r.mCurrIH = CALINK_INVAL;
214 }
215 rowHit = CALINK_INVAL;
216 break;
217 }
218 }
219
220 if (!tParam.Filter(y, z, err2Y, err2Z, GPUCA_MAX_SIN_PHI_LOW)) {
221 rowHit = CALINK_INVAL;
222 break;
223 }
224 CADEBUG(printf("%14s: FIT FILT ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
225 }
226 rowHit = seedIH;
227 r.mNHitsEndRow = ++r.mNHits;
228 r.mLastRow = iRow;
229 r.mEndRow = iRow;
230 r.mNMissed = 0;
231 } while (0);
232
233 /*printf("Extrapolate Row %d X %f Y %f Z %f SinPhi %f DzDs %f QPt %f", iRow, tParam.X(), tParam.Y(), tParam.Z(), tParam.SinPhi(), tParam.DzDs(), tParam.QPt());
234 for (int32_t i = 0;i < 15;i++) printf(" C%d=%6.2f", i, tParam.GetCov(i));
235 printf("\n");*/
236
237 if (r.mCurrIH == CALINK_INVAL) {
238 r.mStage = 1;
239 r.mLastY = tParam.Y(); // Store last spatial position here to start inward following from here
240 r.mLastZ = tParam.Z();
241 if (CAMath::Abs(tParam.SinPhi()) > GPUCA_MAX_SIN_PHI) {
242 r.mGo = 0;
243 }
244 }
245 } else { // forward/backward searching part
246 do {
247 if (r.mStage == 2 && iRow > r.mEndRow) {
248 break;
249 }
250 if (r.mNMissed > tracker.Param().rec.tpc.trackFollowingMaxRowGap) {
251 r.mGo = 0;
252 break;
253 }
254
255 r.mNMissed++;
256
257 float x = row.X();
258 {
259 float tmpY, tmpZ;
260 if (!tParam.GetPropagatedYZ(tracker.Param().bzCLight, x, tmpY, tmpZ)) {
261 r.mGo = 0;
262 rowHit = CALINK_INVAL;
263 break;
264 }
265 tParam.ConstrainZ(tmpZ, tracker.ISector(), z0, r.mLastZ);
266 tracker.GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoX(tracker.ISector(), iRow, tmpY, tmpZ, x);
267 }
268
269 CADEBUG(printf("%14s: SEA TRACK ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
270 if (!tParam.TransportToX(x, tParam.SinPhi(), tParam.GetCosPhi(), tracker.Param().bzCLight, GPUCA_MAX_SIN_PHI_LOW)) {
271 r.mGo = 0;
272 rowHit = CALINK_INVAL;
273 break;
274 }
275 CADEBUG(printf("%14s: SEA PROP ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
276
277 bool found = false;
278 float yUncorrected = tParam.GetY(), zUncorrected = tParam.GetZ();
279 do {
280 if (row.NHits() < 1) {
281 rowHit = CALINK_INVAL;
282 break;
283 }
284
285 GPUglobalref() const cahit2* hits = tracker.HitData(row);
286 GPUglobalref() const calink* firsthit = tracker.FirstHitInBin(row);
287 tracker.GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoNominalYZ(tracker.ISector(), iRow, yUncorrected, zUncorrected, yUncorrected, zUncorrected);
288
289 if (tracker.Param().rec.tpc.rejectEdgeClustersInSeeding && tracker.Param().rejectEdgeClusterByY(yUncorrected, iRow, CAMath::Sqrt(tParam.Err2Y()))) {
290 rowHit = CALINK_INVAL;
291 break;
292 }
293 calink best = CALINK_INVAL;
294
295 float err2Y, err2Z;
296 tracker.GetErrors2Seeding(iRow, *((GPUTPCTrackParam*)&tParam), -1.f, err2Y, err2Z); // TODO: Use correct time
297 if (r.mNHits >= 10) {
298 const float sErr2 = tracker.Param().GetSystematicClusterErrorIFC2(x, tParam.GetY(), tParam.GetZ(), tracker.ISector() >= 18);
299 err2Y += sErr2;
300 err2Z += sErr2;
301 }
302 if (CAMath::Abs(yUncorrected) < x * GPUTPCRow::getTPCMaxY1X()) { // search for the closest hit
303 const float kFactor = tracker.Param().rec.tpc.hitPickUpFactor * tracker.Param().rec.tpc.hitPickUpFactor * 7.0f * 7.0f;
304 const float maxWindow2 = tracker.Param().rec.tpc.hitSearchArea2;
305 const float sy2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Y() + err2Y));
306 const float sz2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Z() + err2Z));
307
308 int32_t bin, ny, nz;
309 row.Grid().GetBinArea(yUncorrected, zUncorrected + tParam.ZOffset(), CAMath::Sqrt(sy2), CAMath::Sqrt(sz2), bin, ny, nz);
310 float ds = 1e6f;
311
312#ifdef __HIPCC__ // Todo: fixme!
313 for (int32_t k = -1; ++k <= nz; /*k++*/) {
314#else
315 for (int32_t k = 0; k <= nz; k++) {
316#endif
317 int32_t nBinsY = row.Grid().Ny();
318 int32_t mybin = bin + k * nBinsY;
319 uint32_t hitFst = firsthit[mybin];
320 uint32_t hitLst = firsthit[mybin + ny + 1];
321#ifdef __HIPCC__ // Todo: fixme!
322 for (uint32_t ih = hitFst - 1; ++ih < hitLst; /*ih++*/) {
323#else
324 for (uint32_t ih = hitFst; ih < hitLst; ih++) {
325#endif
326 cahit2 hh = hits[ih];
327 float y = y0 + hh.x * stepY;
328 float z = z0 + hh.y * stepZ;
329 float dy = y - yUncorrected;
330 float dz = z - zUncorrected;
331 if (dy * dy < sy2 && dz * dz < sz2) {
332 float dds = tracker.Param().rec.tpc.trackFollowingYFactor * CAMath::Abs(dy) + CAMath::Abs(dz);
333 if (dds < ds) {
334 ds = dds;
335 best = ih;
336 }
337 }
338 }
339 }
340 } // end of search for the closest hit
341
342 if (best == CALINK_INVAL) {
343 if (r.mNHits == 0 && r.mStage < 3) {
344 if (rowHit == CALINK_INVAL || rowHit == CALINK_DEAD_CHANNEL) {
345 break;
346 }
347 best = rowHit;
348 } else {
349 rowHit = CALINK_INVAL;
350 break;
351 }
352 }
353
354 cahit2 hh = hits[best];
355 float y = y0 + hh.x * stepY + tParam.GetY() - yUncorrected;
356 float z = z0 + hh.y * stepZ + tParam.GetZ() - zUncorrected;
357
358 CADEBUG(printf("%14s: SEA Hit %5d (%8.3f %8.3f), Res %f %f\n", "", best, y, z, tParam.Y() - y, tParam.Z() - z));
359
360 calink oldHit = (r.mStage == 2 && iRow >= r.mStartRow) ? rowHit : CALINK_INVAL;
361 if (oldHit != best && !tParam.Filter(y, z, err2Y, err2Z, GPUCA_MAX_SIN_PHI_LOW, oldHit != CALINK_INVAL) && r.mNHits != 0) {
362 rowHit = CALINK_INVAL;
363 break;
364 }
365 found = true;
366 rowHit = best;
367 r.mNHits++;
368 r.mNMissed = 0;
369 CADEBUG(printf("%5s hits %3d: SEA FILT ROW %3d X %8.3f -", "", r.mNHits, iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
370 if (r.mStage == 1) {
371 r.mLastRow = iRow;
372 } else {
373 r.mFirstRow = iRow;
374 }
375 } while (false);
376 (void)found;
377 if (!found && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer) {
378 uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(tracker.ISector(), iRow, yUncorrected));
379 if (pad < GPUTPCGeometry::NPads(iRow) && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(tracker.ISector(), iRow, pad)) {
380 r.mNMissed--;
381 rowHit = CALINK_DEAD_CHANNEL;
382 }
383 }
384 } while (0);
385 }
386 if (r.mNHits == 8 && r.mNMissed == 0 && rowHit != CALINK_INVAL && rowHit != CALINK_DEAD_CHANNEL && rowHits && tracker.Param().par.continuousTracking && rowHits[r.mFirstRow] != CALINK_INVAL && rowHits[r.mFirstRow] != CALINK_DEAD_CHANNEL && rowHits[r.mLastRow] != CALINK_INVAL && rowHits[r.mLastRow] != CALINK_DEAD_CHANNEL) {
387 const GPUglobalref() GPUTPCRow& GPUrestrict() row1 = tracker.Row(r.mFirstRow);
388 const GPUglobalref() GPUTPCRow& GPUrestrict() row2 = tracker.Row(r.mLastRow);
389 GPUglobalref() const cahit2* hits1 = tracker.HitData(row1);
390 GPUglobalref() const cahit2* hits2 = tracker.HitData(row2);
391 const cahit2 hh1 = hits1[rowHits[r.mFirstRow]];
392 const cahit2 hh2 = hits2[rowHits[r.mLastRow]];
393 const float z1 = row1.Grid().ZMin() + hh1.y * row1.HstepZ();
394 const float z2 = row2.Grid().ZMin() + hh2.y * row2.HstepZ();
395 float oldOffset = tParam.ZOffset();
396 tParam.ShiftZ(z1, z2, GPUTPCGeometry::Row2X(r.mFirstRow), GPUTPCGeometry::Row2X(r.mLastRow), tracker.Param().bzCLight, tracker.Param().rec.tpc.defaultZOffsetOverR);
397 r.mLastZ -= tParam.ZOffset() - oldOffset;
398 CADEBUG(printf("Shifted z from %f to %f\n", oldOffset, tParam.ZOffset()));
399 }
400}
401
402GPUdic(2, 1) void GPUTPCTrackletConstructor::DoTracklet(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCTrackletConstructor::GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r)
403{
404 int32_t iRow = 0, iRowEnd = GPUCA_ROW_COUNT;
405 GPUTPCTrackParam tParam;
406 calink rowHits[GPUCA_ROW_COUNT];
407 if (r.mGo) {
408 GPUTPCHitId id = tracker.TrackletStartHits()[r.mISH];
409
410 r.mStartRow = r.mEndRow = r.mFirstRow = r.mLastRow = id.RowIndex();
411 r.mCurrIH = id.HitIndex();
412 r.mNMissed = 0;
413 iRow = r.mStartRow;
414 GPUTPCTrackletConstructor::InitTracklet(tParam);
415 }
416 r.mStage = 0;
417 r.mNHits = 0;
418 CADEBUG(printf("Start tracklet\n"));
419
420#ifdef __HIPCC__ // Todo: fixme!
421 for (int32_t iStage = -1; ++iStage < 2; /*iStage++*/) {
422#else
423 for (int32_t iStage = 0; iStage < 2; iStage++) {
424#endif
425 for (; iRow != iRowEnd; iRow += r.mStage == 2 ? -1 : 1) {
426 if (!r.mGo) {
427 break;
428 }
429 UpdateTracklet(0, 0, 0, 0, s, r, tracker, tParam, iRow, rowHits[iRow], rowHits);
430 }
431 if (!r.mGo && r.mStage == 2) {
432 for (; iRow >= r.mStartRow; iRow--) {
433 rowHits[iRow] = CALINK_INVAL;
434 }
435 }
436 if (r.mStage == 2) {
437 StoreTracklet(0, 0, 0, 0, s, r, tracker, tParam, rowHits);
438 } else {
439 r.mStage = 2;
440 r.mNMissed = 0;
441 iRow = r.mEndRow;
442 iRowEnd = -1;
443 float x = tracker.Row(r.mEndRow).X();
444 {
445 float tmpY, tmpZ;
446 if (tParam.GetPropagatedYZ(tracker.Param().bzCLight, x, tmpY, tmpZ)) {
447 if (tracker.ISector() < GPUCA_NSECTORS / 2 ? (tmpZ < 0) : (tmpZ > 0)) {
448 tmpZ = 0;
449 } else if (tracker.ISector() < GPUCA_NSECTORS / 2 ? (tmpZ > GPUTPCGeometry::TPCLength()) : (tmpZ < -GPUTPCGeometry::TPCLength())) {
450 tmpZ = tracker.ISector() < GPUCA_NSECTORS / 2 ? GPUTPCGeometry::TPCLength() : -GPUTPCGeometry::TPCLength();
451 }
452 tracker.GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoX(tracker.ISector(), iRow, tmpY, tmpZ, x);
453 } else {
454 r.mGo = 0;
455 continue;
456 }
457 }
458 if ((r.mGo = (tParam.TransportToX(x, tracker.Param().bzCLight, GPUCA_MAX_SIN_PHI) && tParam.Filter(r.mLastY, r.mLastZ, tParam.Err2Y() * 0.5f, tParam.Err2Z() * 0.5f, GPUCA_MAX_SIN_PHI_LOW, true)))) {
459 CADEBUG(printf("%14s: SEA BACK ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
460 float err2Y, err2Z;
461 tracker.GetErrors2Seeding(r.mEndRow, tParam, -1.f, err2Y, err2Z); // TODO: Use correct time
462 if (tParam.GetCov(0) < err2Y) {
463 tParam.SetCov(0, err2Y);
464 }
465 if (tParam.GetCov(2) < err2Z) {
466 tParam.SetCov(2, err2Z);
467 }
468 CADEBUG(printf("%14s: SEA ADJUS ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n"));
469 r.mNHits -= r.mNHitsEndRow;
470 }
471 }
472 }
473 CADEBUG(printf("End tracklet\n"));
474}
475
476template <>
477GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker)
478{
479 if (get_local_id(0) == 0) {
480 sMem.mNStartHits = *tracker.NStartHits();
481 }
482 CA_SHARED_CACHE(&sMem.mRows[0], tracker.TrackingDataRows(), GPUCA_ROW_COUNT * sizeof(GPUTPCRow));
483 GPUbarrier();
484
485 GPUTPCThreadMemory rMem;
486 for (rMem.mISH = get_global_id(0); rMem.mISH < sMem.mNStartHits; rMem.mISH += get_global_size(0)) {
487 rMem.mGo = 1;
488 DoTracklet(tracker, sMem, rMem);
489 }
490}
491
492template <> // FIXME: GPUgeneric() needed to make the clang spirv output link correctly
493GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking<GPUgeneric() GPUTPCExtrapolationTracking::GPUSharedMemory>(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, GPUTPCTrackParam& GPUrestrict() tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits)
494{
495 GPUTPCThreadMemory rMem;
496 rMem.mISH = iTracklet;
497 rMem.mStage = 3;
498 rMem.mNHits = rMem.mNMissed = 0;
499 rMem.mGo = 1;
500 while (rMem.mGo && row >= 0 && row < GPUCA_ROW_COUNT) {
501 UpdateTracklet(1, 1, 0, 0, sMem, rMem, tracker, tParam, row, rowHits[row], nullptr);
502 row += increment;
503 }
504 if (!CheckCov(tParam)) {
505 rMem.mNHits = 0;
506 }
507 return (rMem.mNHits);
508}
Definition of container class for dE/dx corrections.
Helper class to access correction maps.
int32_t i
#define get_local_id(dim)
#define GPUsharedref()
#define GPUconstantref()
#define GPUdic(...)
#define get_global_size(dim)
#define GPUbarrier()
#define GPUgeneric()
#define GPUrestrict()
#define get_global_id(dim)
#define GPUglobalref()
#define GPUCA_MAX_SIN_PHI_LOW
#define GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(QPTB5)
#define GPUCA_MAX_SIN_PHI
#define CA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared)
Definition GPUDef.h:55
#define CADEBUG(...)
Definition GPUDef.h:79
#define CA_SHARED_CACHE(target, src, size)
Definition GPUDef.h:57
#define CALINK_DEAD_CHANNEL
Definition GPUTPCDef.h:22
#define CALINK_INVAL
Definition GPUTPCDef.h:21
#define GPUCA_NSECTORS
#define GPUCA_ROW_COUNT
GPUdii() void GPUTPCTrackletConstructor
GPUd() bool GPUTPCTrackletConstructor
uint32_t c
Definition RawData.h:2
GLdouble n
Definition glcorearb.h:1982
GLint GLenum GLint x
Definition glcorearb.h:403
GLint y
Definition glcorearb.h:270
typedef void(APIENTRYP PFNGLCULLFACEPROC)(GLenum mode)
GLboolean r
Definition glcorearb.h:1233
GLubyte GLubyte GLubyte GLubyte w
Definition glcorearb.h:852
GLuint GLfloat GLfloat y0
Definition glcorearb.h:5034
GLdouble GLdouble GLdouble z
Definition glcorearb.h:843
uint32_t calink
Definition GPUTPCDef.h:30
o2::mch::DsIndex ds
std::vector< int > row
std::vector< ReadoutWindowData > rows