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) {
135 if (tracker.Param().par.continuousTracking) {
136 tParam.ConstrainZ(z, tracker.ISector(), z0, r.mLastZ);
137 }
138 tracker.GetConstantMem()->calibObjects.fastTransformHelper->TransformXYZ(tracker.ISector(), iRow, x, y, z);
139 }
140 if (iRow == r.mStartRow) {
141 if (tracker.Param().par.continuousTracking) {
142 float refZ = ((z > 0) ? tracker.Param().rec.tpc.defaultZOffsetOverR : -tracker.Param().rec.tpc.defaultZOffsetOverR) * x;
143 float zTmp = refZ;
144 tracker.GetConstantMem()->calibObjects.fastTransformHelper->TransformXYZ(tracker.ISector(), iRow, x, y, zTmp);
145 z += zTmp - refZ; // Add zCorrection (=zTmp - refZ) to z, such that zOffset is set such, that transformed (z - zOffset) becomes refZ
146 tParam.SetZOffset(z - refZ);
147 tParam.SetZ(refZ);
148 r.mLastZ = refZ;
149 } else {
150 tParam.SetZ(z);
151 r.mLastZ = z;
152 tParam.SetZOffset(0.f);
153 }
154 tParam.SetX(x);
155 tParam.SetY(y);
156 r.mLastY = y;
157 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"));
158 } else {
159 float dx = x - tParam.X();
160 float dy, dz;
161 if (r.mNHits >= 10) {
162 dy = y - tParam.Y();
163 dz = z - tParam.Z();
164 } else {
165 dy = y - r.mLastY;
166 dz = z - r.mLastZ;
167 }
168 r.mLastY = y;
169 r.mLastZ = z;
170
171 float ri = 1.f / CAMath::Sqrt(dx * dx + dy * dy);
172 if (iRow == r.mStartRow + 2) {
173 tParam.SetSinPhi(dy * ri);
174 tParam.SetSignCosPhi(dx);
175 tParam.SetDzDs(dz * ri);
176 float err2Y, err2Z;
177 tracker.GetErrors2Seeding(iRow, tParam, -1.f, err2Y, err2Z); // Use correct time
178 tParam.SetCov(0, err2Y);
179 tParam.SetCov(2, err2Z);
180 }
181 float sinPhi, cosPhi;
182 if (r.mNHits >= 10 && CAMath::Abs(tParam.SinPhi()) < GPUCA_MAX_SIN_PHI_LOW) {
183 sinPhi = tParam.SinPhi();
184 cosPhi = CAMath::Sqrt(1 - sinPhi * sinPhi);
185 } else {
186 sinPhi = dy * ri;
187 cosPhi = dx * ri;
188 }
189 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"));
190 if (!tParam.TransportToX(x, sinPhi, cosPhi, tracker.Param().bzCLight, GPUCA_MAX_SIN_PHI)) {
191 rowHit = CALINK_INVAL;
192 break;
193 }
194 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"));
195 float err2Y, err2Z;
196 tracker.GetErrors2Seeding(iRow, tParam.GetZ(), sinPhi, tParam.GetDzDs(), -1.f, err2Y, err2Z); // TODO: Use correct time
197
198 if (r.mNHits >= 10) {
199 const float sErr2 = tracker.Param().GetSystematicClusterErrorIFC2(x, tParam.GetY(), tParam.GetZ(), tracker.ISector() >= 18);
200 err2Y += sErr2;
201 err2Z += sErr2;
202 const float kFactor = tracker.Param().rec.tpc.hitPickUpFactor * tracker.Param().rec.tpc.hitPickUpFactor * 3.5f * 3.5f;
203 float sy2 = kFactor * (tParam.Err2Y() + err2Y);
204 float sz2 = kFactor * (tParam.Err2Z() + err2Z);
205 if (sy2 > tracker.Param().rec.tpc.hitSearchArea2) {
206 sy2 = tracker.Param().rec.tpc.hitSearchArea2;
207 }
208 if (sz2 > tracker.Param().rec.tpc.hitSearchArea2) {
209 sz2 = tracker.Param().rec.tpc.hitSearchArea2;
210 }
211 dy = y - tParam.Y();
212 dz = z - tParam.Z();
213 if (dy * dy > sy2 || dz * dz > sz2) {
214 if (++r.mNMissed >= tracker.Param().rec.tpc.trackFollowingMaxRowGapSeed) {
215 r.mCurrIH = CALINK_INVAL;
216 }
217 rowHit = CALINK_INVAL;
218 break;
219 }
220 }
221
222 if (!tParam.Filter(y, z, err2Y, err2Z, GPUCA_MAX_SIN_PHI_LOW)) {
223 rowHit = CALINK_INVAL;
224 break;
225 }
226 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"));
227 }
228 rowHit = seedIH;
229 r.mNHitsEndRow = ++r.mNHits;
230 r.mLastRow = iRow;
231 r.mEndRow = iRow;
232 r.mNMissed = 0;
233 } while (0);
234
235 /*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());
236 for (int32_t i = 0;i < 15;i++) printf(" C%d=%6.2f", i, tParam.GetCov(i));
237 printf("\n");*/
238
239 if (r.mCurrIH == CALINK_INVAL) {
240 r.mStage = 1;
241 r.mLastY = tParam.Y(); // Store last spatial position here to start inward following from here
242 r.mLastZ = tParam.Z();
243 if (CAMath::Abs(tParam.SinPhi()) > GPUCA_MAX_SIN_PHI) {
244 r.mGo = 0;
245 }
246 }
247 } else { // forward/backward searching part
248 do {
249 if (r.mStage == 2 && iRow > r.mEndRow) {
250 break;
251 }
252 if (r.mNMissed > tracker.Param().rec.tpc.trackFollowingMaxRowGap) {
253 r.mGo = 0;
254 break;
255 }
256
257 r.mNMissed++;
258
259 float x = row.X();
260 {
261 float tmpY, tmpZ;
262 if (!tParam.GetPropagatedYZ(tracker.Param().bzCLight, x, tmpY, tmpZ)) {
263 r.mGo = 0;
264 rowHit = CALINK_INVAL;
265 break;
266 }
267 if (tracker.Param().par.continuousTracking) {
268 tParam.ConstrainZ(tmpZ, tracker.ISector(), z0, r.mLastZ);
269 }
270 tracker.GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoX(tracker.ISector(), iRow, tmpY, tmpZ, x);
271 }
272
273 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"));
274 if (!tParam.TransportToX(x, tParam.SinPhi(), tParam.GetCosPhi(), tracker.Param().bzCLight, GPUCA_MAX_SIN_PHI_LOW)) {
275 r.mGo = 0;
276 rowHit = CALINK_INVAL;
277 break;
278 }
279 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"));
280
281 [[maybe_unused]] bool found = false;
282 float yUncorrected = tParam.GetY(), zUncorrected = tParam.GetZ();
283 do {
284 if (row.NHits() < 1) {
285 rowHit = CALINK_INVAL;
286 break;
287 }
288
289 GPUglobalref() const cahit2* hits = tracker.HitData(row);
290 GPUglobalref() const calink* firsthit = tracker.FirstHitInBin(row);
291 tracker.GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoNominalYZ(tracker.ISector(), iRow, yUncorrected, zUncorrected, yUncorrected, zUncorrected);
292
293 if (tracker.Param().rec.tpc.rejectEdgeClustersInSeeding && tracker.Param().rejectEdgeClusterByY(yUncorrected, iRow, CAMath::Sqrt(tParam.Err2Y()))) {
294 rowHit = CALINK_INVAL;
295 break;
296 }
297 calink best = CALINK_INVAL;
298
299 float err2Y, err2Z;
300 tracker.GetErrors2Seeding(iRow, *((GPUTPCTrackParam*)&tParam), -1.f, err2Y, err2Z); // TODO: Use correct time
301 if (r.mNHits >= 10) {
302 const float sErr2 = tracker.Param().GetSystematicClusterErrorIFC2(x, tParam.GetY(), tParam.GetZ(), tracker.ISector() >= 18);
303 err2Y += sErr2;
304 err2Z += sErr2;
305 }
306 if (CAMath::Abs(yUncorrected) < x * GPUTPCRow::getTPCMaxY1X()) { // search for the closest hit
307 const float kFactor = tracker.Param().rec.tpc.hitPickUpFactor * tracker.Param().rec.tpc.hitPickUpFactor * 7.0f * 7.0f;
308 const float maxWindow2 = tracker.Param().rec.tpc.hitSearchArea2;
309 const float sy2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Y() + err2Y));
310 const float sz2 = CAMath::Min(maxWindow2, kFactor * (tParam.Err2Z() + err2Z));
311
312 int32_t bin, ny, nz;
313 row.Grid().GetBinArea(yUncorrected, zUncorrected + tParam.ZOffset(), CAMath::Sqrt(sy2), CAMath::Sqrt(sz2), bin, ny, nz);
314 float ds = 1e6f;
315
316#ifdef __HIPCC__ // Todo: fixme!
317 for (int32_t k = -1; ++k <= nz; /*k++*/) {
318#else
319 for (int32_t k = 0; k <= nz; k++) {
320#endif
321 int32_t nBinsY = row.Grid().Ny();
322 int32_t mybin = bin + k * nBinsY;
323 uint32_t hitFst = firsthit[mybin];
324 uint32_t hitLst = firsthit[mybin + ny + 1];
325#ifdef __HIPCC__ // Todo: fixme!
326 for (uint32_t ih = hitFst - 1; ++ih < hitLst; /*ih++*/) {
327#else
328 for (uint32_t ih = hitFst; ih < hitLst; ih++) {
329#endif
330 cahit2 hh = hits[ih];
331 float y = y0 + hh.x * stepY;
332 float z = z0 + hh.y * stepZ;
333 float dy = y - yUncorrected;
334 float dz = z - zUncorrected;
335 if (dy * dy < sy2 && dz * dz < sz2) {
336 float dds = tracker.Param().rec.tpc.trackFollowingYFactor * CAMath::Abs(dy) + CAMath::Abs(dz);
337 if (dds < ds) {
338 ds = dds;
339 best = ih;
340 }
341 }
342 }
343 }
344 } // end of search for the closest hit
345
346 if (best == CALINK_INVAL) {
347 if (r.mNHits == 0 && r.mStage < 3) {
348 if (rowHit == CALINK_INVAL || rowHit == CALINK_DEAD_CHANNEL) {
349 break;
350 }
351 best = rowHit;
352 } else {
353 rowHit = CALINK_INVAL;
354 break;
355 }
356 }
357
358 cahit2 hh = hits[best];
359 float y = y0 + hh.x * stepY + tParam.GetY() - yUncorrected;
360 float z = z0 + hh.y * stepZ + tParam.GetZ() - zUncorrected;
361
362 CADEBUG(printf("%14s: SEA Hit %5d (%8.3f %8.3f), Res %f %f\n", "", best, y, z, tParam.Y() - y, tParam.Z() - z));
363
364 calink oldHit = (r.mStage == 2 && iRow >= r.mStartRow) ? rowHit : CALINK_INVAL;
365 if (oldHit != best && !tParam.Filter(y, z, err2Y, err2Z, GPUCA_MAX_SIN_PHI_LOW, oldHit != CALINK_INVAL) && r.mNHits != 0) {
366 rowHit = CALINK_INVAL;
367 break;
368 }
369 found = true;
370 rowHit = best;
371 r.mNHits++;
372 r.mNMissed = 0;
373 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"));
374 if (r.mStage == 1) {
375 r.mLastRow = iRow;
376 } else {
377 r.mFirstRow = iRow;
378 }
379 } while (false);
380 if (!found && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer) {
381 uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(tracker.ISector(), iRow, yUncorrected));
382 if (pad < GPUTPCGeometry::NPads(iRow) && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(tracker.ISector(), iRow, pad)) {
383 r.mNMissed--;
384 rowHit = CALINK_DEAD_CHANNEL;
385 }
386 }
387 } while (0);
388 }
389 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) {
390 const GPUglobalref() GPUTPCRow& GPUrestrict() row1 = tracker.Row(r.mFirstRow);
391 const GPUglobalref() GPUTPCRow& GPUrestrict() row2 = tracker.Row(r.mLastRow);
392 GPUglobalref() const cahit2* hits1 = tracker.HitData(row1);
393 GPUglobalref() const cahit2* hits2 = tracker.HitData(row2);
394 const cahit2 hh1 = hits1[rowHits[r.mFirstRow]];
395 const cahit2 hh2 = hits2[rowHits[r.mLastRow]];
396 const float z1 = row1.Grid().ZMin() + hh1.y * row1.HstepZ();
397 const float z2 = row2.Grid().ZMin() + hh2.y * row2.HstepZ();
398 float oldOffset = tParam.ZOffset();
399 tParam.ShiftZ(z1, z2, GPUTPCGeometry::Row2X(r.mFirstRow), GPUTPCGeometry::Row2X(r.mLastRow), tracker.Param().bzCLight, tracker.Param().rec.tpc.defaultZOffsetOverR);
400 r.mLastZ -= tParam.ZOffset() - oldOffset;
401 CADEBUG(printf("Shifted z from %f to %f\n", oldOffset, tParam.ZOffset()));
402 }
403}
404
405GPUdic(2, 1) void GPUTPCTrackletConstructor::DoTracklet(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCTrackletConstructor::GPUSharedMemory& s, GPUTPCThreadMemory& GPUrestrict() r)
406{
407 int32_t iRow = 0, iRowEnd = GPUCA_ROW_COUNT;
408 GPUTPCTrackParam tParam;
409 calink rowHits[GPUCA_ROW_COUNT];
410 if (r.mGo) {
411 GPUTPCHitId id = tracker.TrackletStartHits()[r.mISH];
412
413 r.mStartRow = r.mEndRow = r.mFirstRow = r.mLastRow = id.RowIndex();
414 r.mCurrIH = id.HitIndex();
415 r.mNMissed = 0;
416 iRow = r.mStartRow;
417 GPUTPCTrackletConstructor::InitTracklet(tParam);
418 }
419 r.mStage = 0;
420 r.mNHits = 0;
421 CADEBUG(printf("Start tracklet\n"));
422
423#ifdef __HIPCC__ // Todo: fixme!
424 for (int32_t iStage = -1; ++iStage < 2; /*iStage++*/) {
425#else
426 for (int32_t iStage = 0; iStage < 2; iStage++) {
427#endif
428 for (; iRow != iRowEnd; iRow += r.mStage == 2 ? -1 : 1) {
429 if (!r.mGo) {
430 break;
431 }
432 UpdateTracklet(0, 0, 0, 0, s, r, tracker, tParam, iRow, rowHits[iRow], rowHits);
433 }
434 if (!r.mGo && r.mStage == 2) {
435 for (; iRow >= r.mStartRow; iRow--) {
436 rowHits[iRow] = CALINK_INVAL;
437 }
438 }
439 if (r.mStage == 2) {
440 StoreTracklet(0, 0, 0, 0, s, r, tracker, tParam, rowHits);
441 } else {
442 r.mStage = 2;
443 r.mNMissed = 0;
444 iRow = r.mEndRow;
445 iRowEnd = -1;
446 float x = tracker.Row(r.mEndRow).X();
447 {
448 float tmpY, tmpZ;
449 if (tParam.GetPropagatedYZ(tracker.Param().bzCLight, x, tmpY, tmpZ)) {
450 if (tracker.ISector() < GPUCA_NSECTORS / 2 ? (tmpZ < 0) : (tmpZ > 0)) {
451 tmpZ = 0;
452 } else if (tracker.ISector() < GPUCA_NSECTORS / 2 ? (tmpZ > GPUTPCGeometry::TPCLength()) : (tmpZ < -GPUTPCGeometry::TPCLength())) {
453 tmpZ = tracker.ISector() < GPUCA_NSECTORS / 2 ? GPUTPCGeometry::TPCLength() : -GPUTPCGeometry::TPCLength();
454 }
455 tracker.GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoX(tracker.ISector(), iRow, tmpY, tmpZ, x);
456 } else {
457 r.mGo = 0;
458 continue;
459 }
460 }
461 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)))) {
462 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"));
463 float err2Y, err2Z;
464 tracker.GetErrors2Seeding(r.mEndRow, tParam, -1.f, err2Y, err2Z); // TODO: Use correct time
465 if (tParam.GetCov(0) < err2Y) {
466 tParam.SetCov(0, err2Y);
467 }
468 if (tParam.GetCov(2) < err2Z) {
469 tParam.SetCov(2, err2Z);
470 }
471 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"));
472 r.mNHits -= r.mNHitsEndRow;
473 }
474 }
475 }
476 CADEBUG(printf("End tracklet\n"));
477}
478
479template <>
480GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker)
481{
482 if (get_local_id(0) == 0) {
483 sMem.mNStartHits = *tracker.NStartHits();
484 }
485 CA_SHARED_CACHE(&sMem.mRows[0], tracker.TrackingDataRows(), GPUCA_ROW_COUNT * sizeof(GPUTPCRow));
486 GPUbarrier();
487
488 GPUTPCThreadMemory rMem;
489 for (rMem.mISH = get_global_id(0); rMem.mISH < sMem.mNStartHits; rMem.mISH += get_global_size(0)) {
490 rMem.mGo = 1;
491 DoTracklet(tracker, sMem, rMem);
492 }
493}
494
495template <> // FIXME: GPUgeneric() needed to make the clang spirv output link correctly
496GPUd() 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)
497{
498 GPUTPCThreadMemory rMem;
499 rMem.mISH = iTracklet;
500 rMem.mStage = 3;
501 rMem.mNHits = rMem.mNMissed = 0;
502 rMem.mGo = 1;
503 while (rMem.mGo && row >= 0 && row < GPUCA_ROW_COUNT) {
504 UpdateTracklet(1, 1, 0, 0, sMem, rMem, tracker, tParam, row, rowHits[row], nullptr);
505 row += increment;
506 }
507 if (!CheckCov(tParam)) {
508 rMem.mNHits = 0;
509 }
510 return (rMem.mNHits);
511}
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:51
#define CADEBUG(...)
Definition GPUDef.h:72
#define CA_SHARED_CACHE(target, src, size)
Definition GPUDef.h:53
#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