15#define GPUCA_CADEBUG 0
16#define GPUCA_MERGE_LOOPER_MC 0
20#if !defined(GPUCA_GPUCODE) && (defined(GPUCA_MERGER_BY_MC_LABEL) || defined(GPUCA_CADEBUG_ENABLED) || GPUCA_MERGE_LOOPER_MC)
25#ifndef GPUCA_GPUCODE_DEVICE
40#include "GPUDefParametersRuntime.h"
68using namespace gputpcgmmergertypes;
70static constexpr int32_t kMaxParts = 400;
96 mNextSectorInd[mid] = 0;
97 mPrevSectorInd[0] = mid;
103#if !defined(GPUCA_GPUCODE) && (defined(GPUCA_MERGER_BY_MC_LABEL) || defined(GPUCA_CADEBUG_ENABLED) || GPUCA_MERGE_LOOPER_MC)
106void GPUTPCGMMerger::CheckMergedTracks()
108 std::vector<bool> trkUsed(SectorTrackInfoLocalTotal());
109 for (int32_t
i = 0;
i < SectorTrackInfoLocalTotal();
i++) {
113 for (int32_t itr = 0; itr < SectorTrackInfoLocalTotal(); itr++) {
115 if (track.PrevSegmentNeighbour() >= 0) {
118 if (track.PrevNeighbour() >= 0) {
124 int32_t iTrk = tr - mSectorTrackInfos;
126 GPUError(
"FAILURE: double use");
128 trkUsed[iTrk] =
true;
130 int32_t jtr = tr->NextSegmentNeighbour();
132 tr = &(mSectorTrackInfos[jtr]);
135 jtr = trbase->NextNeighbour();
137 trbase = &(mSectorTrackInfos[jtr]);
139 if (tr->PrevSegmentNeighbour() >= 0) {
148 for (int32_t
i = 0;
i < SectorTrackInfoLocalTotal();
i++) {
149 if (trkUsed[
i] ==
false) {
150 GPUError(
"FAILURE: trk missed");
166template <
class T,
class S>
167int64_t GPUTPCGMMerger::GetTrackLabelA(
const S& trk)
const
171 if constexpr (std::is_same<S, GPUTPCGMBorderTrack&>::value) {
172 sectorTrack = &mSectorTrackInfos[trk.TrackID()];
173 nClusters = sectorTrack->OrigTrack()->NHits();
177 auto acc =
GPUTPCTrkLbl<false, GPUTPCTrkLbl_ret>(resolveMCLabels<T>(GetConstantMem()->ioPtrs.clustersNative ? GetConstantMem()->ioPtrs.clustersNative->clustersMCTruth : nullptr, GetConstantMem()->ioPtrs.mcLabelsTPC), 0.5f);
180 if constexpr (std::is_same<S, GPUTPCGMBorderTrack&>::value) {
181 const GPUTPCTracker& tracker = GetConstantMem()->tpcTrackers[sectorTrack->Sector()];
182 const GPUTPCHitId& ic = tracker.TrackHits()[sectorTrack->OrigTrack()->FirstHitID() +
i];
183 id = tracker.Data().ClusterDataIndex(tracker.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sectorTrack->Sector()][0];
185 id = mClusters[trk.FirstClusterRef() +
i].
num;
189 return acc.computeLabel().id;
193int64_t GPUTPCGMMerger::GetTrackLabel(
const S& trk)
const
195#ifdef GPUCA_TPC_GEOMETRY_O2
196 if (GetConstantMem()->ioPtrs.clustersNative->clustersMCTruth) {
197 return GetTrackLabelA<o2::dataformats::ConstMCTruthContainerView<o2::MCCompLabel>,
S>(trk);
201 return GetTrackLabelA<AliHLTTPCClusterMCLabel, S>(trk);
208void GPUTPCGMMerger::PrintMergeGraph(
const GPUTPCGMSectorTrack* trk, std::ostream& out)
const
211 while (trk->PrevSegmentNeighbour() >= 0) {
212 trk = &mSectorTrackInfos[trk->PrevSegmentNeighbour()];
215 while (trk->PrevNeighbour() >= 0) {
216 trk = &mSectorTrackInfos[trk->PrevNeighbour()];
219 int32_t nextId = trk - mSectorTrackInfos;
220 out <<
"Graph of track " << (orgTrack - mSectorTrackInfos) <<
"\n";
221 while (nextId >= 0) {
222 trk = &mSectorTrackInfos[nextId];
223 if (trk->PrevSegmentNeighbour() >= 0) {
224 out <<
"TRACK TREE INVALID!!! " << trk->PrevSegmentNeighbour() <<
" --> " << nextId <<
"\n";
226 out << (trk == orgTower ?
"--" :
" ");
227 while (nextId >= 0) {
229 if (trk != trk2 && (trk2->PrevNeighbour() >= 0 || trk2->NextNeighbour() >= 0)) {
230 out <<
" (TRACK TREE INVALID!!! " << trk2->PrevNeighbour() <<
" <-- " << nextId <<
" --> " << trk2->NextNeighbour() <<
") ";
233 snprintf(tmp, 128,
" %s%5d(%5.2f)", trk2 == orgTrack ?
"!" :
" ", nextId, trk2->QPt());
235 nextId = trk2->NextSegmentNeighbour();
238 nextId = trk->NextNeighbour();
258 mBorder[
iSector] = mBorderMemory + 2 * nTracks;
260 mBorderRange[
iSector] = mBorderRangeMemory + 2 * nTracks;
268 memMax = (
void*)std::max((
size_t)mem, (size_t)memMax);
272 memMax = (
void*)std::max((
size_t)mem, (size_t)memMax);
276 memMax = (
void*)std::max((
size_t)mem, (size_t)memMax);
279 memMax = (
void*)std::max((
size_t)mem, (size_t)memMax);
293 if (mergerSortTracks) {
319 mClusterStateExt =
nullptr;
369 mNTotalSectorTracks = 0;
371 mNMaxSingleSectorTracks = 0;
374 mNTotalSectorTracks += ntrk;
376 if (mNMaxSingleSectorTracks < ntrk) {
377 mNMaxSingleSectorTracks = ntrk;
381 if (CAMath::Abs(
Param().polynomialField.GetNominalBz()) < (gpu_common_constants::kZeroFieldCut * gpu_common_constants::kCLight)) {
394 mNMaxClusters = mNClusters;
396 mNMaxLooperMatches = mNMaxClusters / 4;
403 throw std::runtime_error(
"mNMaxSingleSectorTracks too small");
407 throw std::runtime_error(
"Must run also sector tracking");
416 const int32_t
n =
output ? mMemory->nOutputTracks : SectorTrackInfoLocalTotal();
417 for (int32_t
i = iBlock * nThreads + iThread;
i <
n;
i += nThreads * nBlocks) {
425 prop.SetMaterialTPC();
427 prop.SetToyMCEventsFlag(
false);
428 prop.SetSeedingErrors(
true);
429 prop.SetFitInProjections(
false);
430 prop.SetPolynomialField(&
Param().polynomialField);
432 trk.X() = inTrack->Param().GetX();
433 trk.
Y() = inTrack->Param().GetY();
434 trk.
Z() = inTrack->Param().GetZ();
435 trk.SinPhi() = inTrack->Param().GetSinPhi();
436 trk.DzDs() = inTrack->Param().GetDzDs();
437 trk.QPt() = inTrack->Param().GetQPt();
438 trk.TZOffset() =
Param().par.earlyTpcTransform ? inTrack->Param().GetZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, inTrack->Param().GetZOffset(),
Param().continuousMaxTimeBin);
439 trk.ShiftZ(
this, sector, sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN(), inTrack->Param().GetX(), inTrack->Param().GetX());
440 sectorTrack.SetX2(0.f);
441 for (int32_t way = 0; way < 2; way++) {
443 prop.SetFitInProjections(
true);
444 prop.SetPropagateBzOnly(
true);
446 trk.ResetCovariance();
447 prop.SetTrack(&trk,
alpha);
448 int32_t
start = way ? inTrack->NHits() - 1 : 0;
449 int32_t
end = way ? 0 : (inTrack->NHits() - 1);
450 int32_t incr = way ? -1 : 1;
454 const GPUTPCTracker& tracker = GetConstantMem()->tpcTrackers[sector];
455 const GPUTPCHitId& ic = tracker.TrackHits()[inTrack->FirstHitID() +
i];
456 int32_t clusterIndex = tracker.Data().ClusterDataIndex(tracker.Data().Row(ic.RowIndex()), ic.HitIndex());
458 const ClusterNative& cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sector][0] + clusterIndex];
459 flags = cl.getFlags();
460 if (
Param().
par.earlyTpcTransform) {
461 x = tracker.Data().ClusterData()[clusterIndex].x;
462 y = tracker.Data().ClusterData()[clusterIndex].y;
463 z = tracker.Data().ClusterData()[clusterIndex].
z - trk.TZOffset();
465 GetConstantMem()->calibObjects.fastTransformHelper->Transform(sector,
row, cl.getPad(), cl.getTime(),
x,
y,
z, trk.TZOffset());
467 if (prop.PropagateToXAlpha(
x,
alpha,
true)) {
470 trk.ConstrainSinPhi();
471 if (prop.Update(
y,
z,
row,
Param(),
flags &
GPUTPCGMMergedTrackHit::clustererAndSharedFlags, 0,
nullptr,
false, sector, -1.f, 0.f, 0.f)) {
474 trk.ConstrainSinPhi();
477 sectorTrack.SetParam2(trk);
479 sectorTrack.Set(trk, inTrack,
alpha, sector);
487 const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
488 const GPUTPCHitId& ic1 = trk.TrackHits()[sectorTr->FirstHitID()];
489 const GPUTPCHitId& ic2 = trk.TrackHits()[sectorTr->FirstHitID() + sectorTr->NHits() - 1];
490 int32_t clusterIndex1 = trk.Data().ClusterDataIndex(trk.Data().Row(ic1.RowIndex()), ic1.HitIndex());
491 int32_t clusterIndex2 = trk.Data().ClusterDataIndex(trk.Data().Row(ic2.RowIndex()), ic2.HitIndex());
492 if (
Param().
par.earlyTpcTransform) {
493 track.SetClusterZT(trk.Data().ClusterData()[clusterIndex1].
z, trk.Data().ClusterData()[clusterIndex2].
z);
495 const ClusterNative* cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[iSector][0];
496 track.SetClusterZT(cl[clusterIndex1].
getTime(), cl[clusterIndex2].
getTime());
502 mSectorTrackInfoIndex[
id] = mMemory->nUnpackedTracks;
505GPUd()
void GPUTPCGMMerger::UnpackSectorGlobal(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector)
507 const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
509 const GPUTPCTrack* sectorTr = mMemory->firstExtrapolatedTracks[iSector];
510 uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks;
511 uint32_t nTracks = *trk.NTracks();
512 for (uint32_t itr = nLocalTracks + iBlock * nThreads + iThread; itr < nTracks; itr += nBlocks * nThreads) {
513 sectorTr = &trk.Tracks()[itr];
514 int32_t localId = mTrackIDs[(sectorTr->LocalTrackId() >> 24) * mNMaxSingleSectorTracks + (sectorTr->LocalTrackId() & 0xFFFFFF)];
518 uint32_t myTrack = CAMath::AtomicAdd(&mMemory->nUnpackedTracks, 1u);
520 SetTrackClusterZT(track, iSector, sectorTr);
521 track.Set(
this, sectorTr,
alpha, iSector);
522 track.SetGlobalSectorTrackCov();
523 track.SetPrevNeighbour(-1);
524 track.SetNextNeighbour(-1);
525 track.SetNextSegmentNeighbour(-1);
526 track.SetPrevSegmentNeighbour(-1);
527 track.SetLocalTrackId(localId);
531GPUd()
void GPUTPCGMMerger::UnpackResetIds(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector)
533 const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
534 uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks;
535 for (uint32_t
i = iBlock * nThreads + iThread;
i < nLocalTracks;
i += nBlocks * nThreads) {
536 mTrackIDs[iSector * mNMaxSingleSectorTracks +
i] = -1;
540GPUd()
void GPUTPCGMMerger::RefitSectorTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector)
542 const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSector];
543 uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks;
548 for (uint32_t itr = iBlock * nThreads + iThread; itr < nLocalTracks; itr += nBlocks * nThreads) {
549 sectorTr = &trk.Tracks()[itr];
551 SetTrackClusterZT(track, iSector, sectorTr);
552 if (
Param().
rec.tpc.mergerCovSource == 0) {
553 track.Set(
this, sectorTr,
alpha, iSector);
557 }
else if (
Param().
rec.tpc.mergerCovSource == 1) {
558 track.Set(
this, sectorTr,
alpha, iSector);
559 track.CopyBaseTrackCov();
560 }
else if (
Param().
rec.tpc.mergerCovSource == 2) {
561 if (RefitSectorTrack(track, sectorTr,
alpha, iSector)) {
562 track.Set(
this, sectorTr,
alpha, iSector);
569 CADEBUG(GPUInfo(
"INPUT Sector %d, Track %u, QPt %f DzDs %f", iSector, itr, track.QPt(), track.DzDs()));
570 track.SetPrevNeighbour(-1);
571 track.SetNextNeighbour(-1);
572 track.SetNextSegmentNeighbour(-1);
573 track.SetPrevSegmentNeighbour(-1);
574 track.SetExtrapolatedTrackId(0, -1);
575 track.SetExtrapolatedTrackId(1, -1);
576 uint32_t myTrack = CAMath::AtomicAdd(&mMemory->nUnpackedTracks, 1u);
577 mTrackIDs[iSector * mNMaxSingleSectorTracks + sectorTr->LocalTrackId()] = myTrack;
578 mSectorTrackInfos[myTrack] = track;
582GPUd()
void GPUTPCGMMerger::LinkExtrapolatedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
584 for (int32_t itr = SectorTrackInfoGlobalFirst(0) + iBlock * nThreads + iThread; itr < SectorTrackInfoGlobalLast(NSECTORS - 1); itr += nThreads * nBlocks) {
587 if (localTrack.ExtrapolatedTrackId(0) != -1 || !CAMath::AtomicCAS(&localTrack.ExtrapolatedTrackIds()[0], -1, itr)) {
588 localTrack.SetExtrapolatedTrackId(1, itr);
598 float fieldBz =
Param().bzCLight;
600 float dAlpha =
Param().par.dAlpha / 2;
604 dAlpha = dAlpha - CAMath::Pi() / 2;
605 }
else if (iBorder == 1) {
606 dAlpha = -dAlpha - CAMath::Pi() / 2;
607 }
else if (iBorder == 2) {
608 x0 = GPUTPCGeometry::Row2X(63);
609 }
else if (iBorder == 3) {
611 x0 = GPUTPCGeometry::Row2X(63);
612 }
else if (iBorder == 4) {
614 x0 = GPUTPCGeometry::Row2X(63);
617 const float maxSin = CAMath::Sin(60.f / 180.f * CAMath::Pi());
618 float cosAlpha = CAMath::Cos(dAlpha);
619 float sinAlpha = CAMath::Sin(dAlpha);
622 for (int32_t itr = iBlock * nThreads + iThread; itr < SectorTrackInfoLocalTotal(); itr += nThreads * nBlocks) {
624 int32_t iSector = track->Sector();
626 if (track->PrevSegmentNeighbour() >= 0 && track->Sector() == mSectorTrackInfos[track->PrevSegmentNeighbour()].Sector()) {
629 if (useOrigTrackParam) {
630 if (CAMath::Abs(track->QPt()) *
Param().qptB5Scaler <
Param().rec.tpc.mergerLooperQPtB5Limit) {
634 while (track->NextSegmentNeighbour() >= 0 && track->Sector() == mSectorTrackInfos[track->NextSegmentNeighbour()].Sector()) {
635 track = &mSectorTrackInfos[track->NextSegmentNeighbour()];
636 if (track->OrigTrack()->Param().X() < trackMin->OrigTrack()->Param().X()) {
640 trackTmp = *trackMin;
642 if (
Param().rec.tpc.mergerCovSource == 2 && trackTmp.X2() != 0.f) {
643 trackTmp.UseParam2();
645 trackTmp.Set(
this, trackMin->OrigTrack(), trackMin->Alpha(), trackMin->Sector());
648 if (CAMath::Abs(track->QPt()) *
Param().qptB5Scaler <
Param().rec.tpc.mergerLooperSecondHorizontalQPtB5Limit) {
649 if (iBorder == 0 && track->NextNeighbour() >= 0) {
652 if (iBorder == 1 && track->PrevNeighbour() >= 0) {
659 if (track->TransportToXAlpha(
this,
x0, sinAlpha, cosAlpha, fieldBz,
b, maxSin)) {
661 b.SetNClusters(track->NClusters());
662 for (int32_t
i = 0;
i < 4;
i++) {
663 if (CAMath::Abs(
b.Cov()[
i]) >= 5.0f) {
667 if (CAMath::Abs(
b.Cov()[4]) >= 0.5f) {
670 uint32_t myTrack = CAMath::AtomicAdd(&nB[iSector], 1u);
671 B[iSector][myTrack] =
b;
677GPUd()
void GPUTPCGMMerger::MergeBorderTracks<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector1,
GPUTPCGMBorderTrack*
B1, int32_t N1, int32_t iSector2,
GPUTPCGMBorderTrack* B2, int32_t N2, int32_t mergeMode)
679 CADEBUG(GPUInfo(
"\nMERGING Sectors %d %d NTracks %d %d CROSS %d", iSector1, iSector2, N1, N2, mergeMode));
681 GPUTPCGMBorderRange* range2 = mBorderRange[iSector2] + *GetConstantMem()->tpcTrackers[iSector2].NTracks();
682 bool sameSector = (iSector1 == iSector2);
683 for (int32_t itr = iBlock * nThreads + iThread; itr < N1; itr += nThreads * nBlocks) {
685 float d = CAMath::Max(0.5f, 3.5f * CAMath::Sqrt(
b.Cov()[1]));
686 if (CAMath::Abs(
b.Par()[4]) *
Param().qptB5Scaler >= 20) {
691 CADEBUG(printf(
" Input Sector 1 %d Track %d: ", iSector1, itr);
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b.Par()[
i]); } printf(
" - ");
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b.Cov()[
i]); } printf(
" - D %8.3f\n", d));
694 range.fMin =
b.Par()[1] +
b.ZOffsetLinear() - d;
695 range.fMax =
b.Par()[1] +
b.ZOffsetLinear() + d;
702 for (int32_t itr = iBlock * nThreads + iThread; itr < N2; itr += nThreads * nBlocks) {
704 float d = CAMath::Max(0.5f, 3.5f * CAMath::Sqrt(
b.Cov()[1]));
705 if (CAMath::Abs(
b.Par()[4]) *
Param().qptB5Scaler >= 20) {
710 CADEBUG(printf(
" Input Sector 2 %d Track %d: ", iSector2, itr);
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b.Par()[
i]); } printf(
" - ");
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b.Cov()[
i]); } printf(
" - D %8.3f\n", d));
713 range.fMin =
b.Par()[1] +
b.ZOffsetLinear() - d;
714 range.fMax =
b.Par()[1] +
b.ZOffsetLinear() + d;
721GPUd()
void GPUTPCGMMerger::MergeBorderTracks<1>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector1,
GPUTPCGMBorderTrack*
B1, int32_t N1, int32_t iSector2,
GPUTPCGMBorderTrack* B2, int32_t N2, int32_t mergeMode)
723#if !defined(GPUCA_GPUCODE_COMPILEKERNELS)
725 GPUTPCGMBorderRange* range2 = mBorderRange[iSector2] + *GetConstantMem()->tpcTrackers[iSector2].NTracks();
730 }
else if (iBlock == 1) {
735 printf(
"This sorting variant is disabled for RTC");
742#ifndef GPUCA_SPECIALIZE_THRUST_SORTS
754GPUd()
void GPUTPCGMMerger::MergeBorderTracks<2>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector1,
GPUTPCGMBorderTrack*
B1, int32_t N1, int32_t iSector2,
GPUTPCGMBorderTrack* B2, int32_t N2, int32_t mergeMode)
757 float factor2ys =
Param().rec.tpc.trackMergerFactor2YS;
758 float factor2zt =
Param().rec.tpc.trackMergerFactor2ZT;
759 float factor2k =
Param().rec.tpc.trackMergerFactor2K;
760 float factor2General =
Param().rec.tpc.trackMergerFactor2General;
762 factor2k = factor2General * factor2k;
763 factor2ys = factor2General * factor2ys;
764 factor2zt = factor2General * factor2zt;
766 int32_t minNPartHits =
Param().rec.tpc.trackMergerMinPartHits;
767 int32_t minNTotalHits =
Param().rec.tpc.trackMergerMinTotalHits;
769 bool sameSector = (iSector1 == iSector2);
772 GPUTPCGMBorderRange* range2 = mBorderRange[iSector2] + *GetConstantMem()->tpcTrackers[iSector2].NTracks();
775 for (int32_t i1 = iBlock * nThreads + iThread; i1 < N1; i1 += nThreads * nBlocks) {
777 while (i2 < N2 && range2[i2].fMax < r1.
fMin) {
782 if (
b1.NClusters() < minNPartHits) {
788 for (int32_t k2 = i2; k2 < N2; k2++) {
793 if (sameSector && (r1.
fId >= r2.
fId)) {
799#if defined(GPUCA_MERGER_BY_MC_LABEL) && !defined(GPUCA_GPUCODE)
800 int64_t label1 = GetTrackLabel(
b1);
801 int64_t label2 = GetTrackLabel(b2);
802 if (label1 != label2 && label1 != -1)
805 CADEBUG(
if (GetConstantMem()->ioPtrs.mcLabelsTPC) {printf(
"Comparing track %3d to %3d: ", r1.fId, r2.fId); for (int32_t i = 0; i < 5; i++) { printf(
"%8.3f ", b1.Par()[i]); } printf(
" - ");
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b1.Cov()[
i]); } printf(
"\n%28s",
""); });
806 CADEBUG(
if (GetConstantMem()->ioPtrs.mcLabelsTPC) {for (int32_t i = 0; i < 5; i++) { printf(
"%8.3f ", b2.Par()[i]); } printf(
" - ");
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ", b2.Cov()[
i]); } printf(
" - %5s - ", GetTrackLabel(
b1) == GetTrackLabel(b2) ?
"CLONE" :
"FAKE"); });
807 if (b2.NClusters() < lBest2) {
808 CADEBUG2(
continue, printf(
"!NCl1\n"));
812 int32_t maxRowDiff = mergeMode == 2 ? 1 : 3;
813 if (CAMath::Abs(
b1.Row() - b2.Row()) > maxRowDiff) {
814 CADEBUG2(
continue, printf(
"!ROW\n"));
816 if (CAMath::Abs(
b1.Par()[2] - b2.Par()[2]) > 0.5f || CAMath::Abs(
b1.Par()[3] - b2.Par()[3]) > 0.5f) {
817 CADEBUG2(
continue, printf(
"!CE SinPhi/Tgl\n"));
823 if (!
b1.CheckChi2Y(b2, factor2ys)) {
827 if (!
b1.CheckChi2QPt(b2, factor2k)) {
828 CADEBUG2(
continue, printf(
"!QPt\n"));
830 float fys = CAMath::Abs(
b1.Par()[4]) *
Param().qptB5Scaler < 20 ? factor2ys : (2.f * factor2ys);
831 float fzt = CAMath::Abs(
b1.Par()[4]) *
Param().qptB5Scaler < 20 ? factor2zt : (2.f * factor2zt);
832 if (!
b1.CheckChi2YS(b2, fys)) {
833 CADEBUG2(
continue, printf(
"!YS\n"));
835 if (!
b1.CheckChi2ZT(b2, fzt)) {
836 CADEBUG2(
continue, printf(
"!ZT\n"));
838 if (CAMath::Abs(
b1.Par()[4]) *
Param().qptB5Scaler < 20) {
839 if (b2.NClusters() < minNPartHits) {
840 CADEBUG2(
continue, printf(
"!NCl2\n"));
842 if (
b1.NClusters() + b2.NClusters() < minNTotalHits) {
843 CADEBUG2(
continue, printf(
"!NCl3\n"));
846 CADEBUG(printf(
"OK: dZ %8.3f D1 %8.3f D2 %8.3f\n", CAMath::Abs(
b1.Par()[1] - b2.Par()[1]), 3.5f * sqrt(
b1.Cov()[1]), 3.5f * sqrt(b2.Cov()[1])));
848 lBest2 = b2.NClusters();
849 iBest2 = b2.TrackID();
855 GPUCA_DEBUG_STREAMER_CHECK(
float weight =
b1.Par()[4] *
b1.Par()[4];
if (o2::utils::DebugStreamer::checkStream(
o2::utils::StreamFlags::streamMergeBorderTracksBest,
b1.TrackID(),
weight)) { MergedTrackStreamer(
b1, MergedTrackStreamerFindBorderTrack(B2, N2, iBest2),
"merge_best_track", iSector1, iSector2, mergeMode,
weight, o2::utils::DebugStreamer::getSamplingFrequency(
o2::utils::StreamFlags::streamMergeBorderTracksBest)); });
859 CADEBUG(GPUInfo(
"Found match %d %d",
b1.TrackID(), iBest2));
861 mTrackLinks[
b1.TrackID()] = iBest2;
871 if (withinSector == 1) {
873 n1 = n2 = mMemory->tmpCounter[iSector];
874 b1 = b2 = mBorder[iSector];
875 }
else if (withinSector == -1) {
876 jSector = (iSector + NSECTORS / 2);
877 const int32_t
offset = mergeMode == 2 ? NSECTORS : 0;
878 n1 = mMemory->tmpCounter[iSector +
offset];
879 n2 = mMemory->tmpCounter[jSector +
offset];
881 b2 = mBorder[jSector +
offset];
883 jSector = mNextSectorInd[iSector];
884 n1 = mMemory->tmpCounter[iSector];
885 n2 = mMemory->tmpCounter[NSECTORS + jSector];
886 b1 = mBorder[iSector];
887 b2 = mBorder[NSECTORS + jSector];
892GPUd()
void GPUTPCGMMerger::MergeBorderTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector, int8_t withinSector, int8_t mergeMode)
897 MergeBorderTracksSetup(n1, n2,
b1, b2, jSector, iSector, withinSector, mergeMode);
898 MergeBorderTracks<I>(nBlocks, nThreads, iBlock, iThread, iSector,
b1, n1, jSector, b2, n2, mergeMode);
901#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE)
902template GPUdni()
void GPUTPCGMMerger::MergeBorderTracks<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector, int8_t withinSector, int8_t mergeMode);
903template
GPUdni()
void GPUTPCGMMerger::MergeBorderTracks<1>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector, int8_t withinSector, int8_t mergeMode);
904template
GPUdni()
void GPUTPCGMMerger::MergeBorderTracks<2>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t iSector, int8_t withinSector, int8_t mergeMode);
907GPUd()
void GPUTPCGMMerger::MergeWithinSectorsPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
909 float x0 = GPUTPCGeometry::Row2X(63);
910 const float maxSin = CAMath::Sin(60.f / 180.f * CAMath::Pi());
912 for (int32_t itr = iBlock * nThreads + iThread; itr < SectorTrackInfoLocalTotal(); itr += nThreads * nBlocks) {
914 int32_t iSector = track.Sector();
916 if (track.TransportToX(
this,
x0,
Param().bzCLight,
b, maxSin)) {
918 CADEBUG(printf(
"WITHIN SECTOR %d Track %d - ", iSector, itr);
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b.Par()[
i]); } printf(
" - ");
for (int32_t
i = 0;
i < 5;
i++) { printf(
"%8.3f ",
b.Cov()[
i]); } printf(
"\n"));
919 b.SetNClusters(track.NClusters());
920 uint32_t myTrack = CAMath::AtomicAdd(&mMemory->tmpCounter[iSector], 1u);
921 mBorder[iSector][myTrack] =
b;
926GPUd()
void GPUTPCGMMerger::MergeSectorsPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, int32_t border0, int32_t border1, int8_t useOrigTrackParam)
928 bool part2 = iBlock & 1;
929 int32_t
border = part2 ? border1 : border0;
936 MergeSectorsPrepareStep2((nBlocks + !part2) >> 1, nThreads, iBlock >> 1, iThread,
border,
b,
n, useOrigTrackParam);
941 start = (elems + nBlocks - 1) / nBlocks * iBlock;
942 end = (elems + nBlocks - 1) / nBlocks * (iBlock + 1);
943 end = CAMath::Min(elems,
end);
952 u = mTrackCCRoots[u];
953 v = mTrackCCRoots[
v];
957 int32_t
h = CAMath::Max(u,
v);
958 int32_t l = CAMath::Min(u,
v);
960 int32_t old = CAMath::AtomicCAS(&mTrackCCRoots[
h],
h, l);
965 u = mTrackCCRoots[
h];
970GPUd()
void GPUTPCGMMerger::ResolveFindConnectedComponentsSetup(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
973 setBlockRange(SectorTrackInfoLocalTotal(), nBlocks, iBlock,
start,
end);
974 for (int32_t
i =
start + iThread;
i <
end;
i += nThreads) {
975 mTrackCCRoots[
i] =
i;
979GPUd()
void GPUTPCGMMerger::ResolveFindConnectedComponentsHookLinks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
984 setBlockRange(SectorTrackInfoLocalTotal(), nBlocks, iBlock,
start,
end);
985 for (int32_t itr =
start + iThread; itr <
end; itr += nThreads) {
986 hookEdge(itr, mTrackLinks[itr]);
990GPUd()
void GPUTPCGMMerger::ResolveFindConnectedComponentsHookNeighbors(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
993 nBlocks = nBlocks / 4 * 4;
994 if (iBlock >= nBlocks) {
999 setBlockRange(SectorTrackInfoLocalTotal(), nBlocks / 4, iBlock / 4,
start,
end);
1001 int32_t myNeighbor = iBlock % 4;
1003 for (int32_t itr =
start + iThread; itr <
end; itr += nThreads) {
1004 int32_t
v = mSectorTrackInfos[itr].AnyNeighbour(myNeighbor);
1009GPUd()
void GPUTPCGMMerger::ResolveFindConnectedComponentsMultiJump(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1013 setBlockRange(SectorTrackInfoLocalTotal(), nBlocks, iBlock,
start,
end);
1014 for (int32_t itr =
start + iThread; itr <
end; itr += nThreads) {
1016 int32_t next = mTrackCCRoots[root];
1022 next = mTrackCCRoots[next];
1023 }
while (root != next);
1024 mTrackCCRoots[itr] = root;
1061 setBlockRange(SectorTrackInfoLocalTotal(), nBlocks, iBlock,
start,
end);
1063 for (int32_t baseIdx = 0; baseIdx < SectorTrackInfoLocalTotal(); baseIdx += nThreads) {
1064 int32_t itr = baseIdx + iThread;
1065 bool inRange = itr < SectorTrackInfoLocalTotal();
1069 itr2 = mTrackLinks[itr];
1072 bool resolveSector = (itr2 > -1);
1073 if (resolveSector) {
1074 int32_t root = mTrackCCRoots[itr];
1075 resolveSector &= (
start <= root) && (root <
end);
1078 int16_t smemIdx = work_group_scan_inclusive_add(int16_t(resolveSector));
1080 if (resolveSector) {
1081 smem.iTrack1[smemIdx - 1] = itr;
1082 smem.iTrack2[smemIdx - 1] = itr2;
1086 if (iThread < nThreads - 1) {
1090 const int32_t nSectors = smemIdx;
1092 for (int32_t
i = 0;
i < nSectors;
i++) {
1093 itr = smem.iTrack1[
i];
1094 itr2 = smem.iTrack2[
i];
1101 bool sameSegment = CAMath::Abs(track1->NClusters() > track2->NClusters() ? track1->QPt() : track2->QPt()) *
Param().qptB5Scaler < 2 || track1->QPt() * track2->QPt() > 0;
1106 while (track2->PrevSegmentNeighbour() >= 0) {
1107 track2 = &mSectorTrackInfos[track2->PrevSegmentNeighbour()];
1110 if (track1 == track2) {
1113 while (track1->PrevSegmentNeighbour() >= 0) {
1114 track1 = &mSectorTrackInfos[track1->PrevSegmentNeighbour()];
1115 if (track1 == track2) {
1119 GPUCommonAlgorithm::swap(track1, track1Base);
1120 for (int32_t k = 0; k < 2; k++) {
1122 while (tmp->Neighbour(k) >= 0) {
1123 tmp = &mSectorTrackInfos[tmp->Neighbour(k)];
1124 if (tmp == track2) {
1130 while (track1->NextSegmentNeighbour() >= 0) {
1131 track1 = &mSectorTrackInfos[track1->NextSegmentNeighbour()];
1132 if (track1 == track2) {
1137 while (track1->PrevSegmentNeighbour() >= 0) {
1138 track1 = &mSectorTrackInfos[track1->PrevSegmentNeighbour()];
1141 if (track1 == track2) {
1144 for (int32_t k = 0; k < 2; k++) {
1146 while (tmp->Neighbour(k) >= 0) {
1147 tmp = &mSectorTrackInfos[tmp->Neighbour(k)];
1148 if (tmp == track2) {
1154 float z1min, z1max, z2min, z2max;
1155 z1min = track1->MinClusterZT();
1156 z1max = track1->MaxClusterZT();
1157 z2min = track2->MinClusterZT();
1158 z2max = track2->MaxClusterZT();
1159 if (track1 != track1Base) {
1160 z1min = CAMath::Min(z1min, track1Base->MinClusterZT());
1161 z1max = CAMath::Max(z1max, track1Base->MaxClusterZT());
1163 if (track2 != track2Base) {
1164 z2min = CAMath::Min(z2min, track2Base->MinClusterZT());
1165 z2max = CAMath::Max(z2max, track2Base->MaxClusterZT());
1167 bool goUp = z2max - z1min > z1max - z2min;
1169 if (track1->Neighbour(goUp) < 0 && track2->Neighbour(!goUp) < 0) {
1170 track1->SetNeighbor(track2 - mSectorTrackInfos, goUp);
1171 track2->SetNeighbor(track1 - mSectorTrackInfos, !goUp);
1175 }
else if (track1->Neighbour(goUp) < 0) {
1176 track2 = &mSectorTrackInfos[track2->Neighbour(!goUp)];
1177 GPUCommonAlgorithm::swap(track1, track2);
1178 }
else if (track2->Neighbour(!goUp) < 0) {
1179 track1 = &mSectorTrackInfos[track1->Neighbour(goUp)];
1181 track1 = &mSectorTrackInfos[track1->Neighbour(goUp)];
1183 track1Base = track1;
1186 track2Base = track2;
1188 while (track1->NextSegmentNeighbour() >= 0) {
1189 track1 = &mSectorTrackInfos[track1->NextSegmentNeighbour()];
1192 track1->SetNextSegmentNeighbour(track2 - mSectorTrackInfos);
1193 track2->SetPrevSegmentNeighbour(track1 - mSectorTrackInfos);
1196 for (int32_t k = 0; k < 2; k++) {
1197 track1 = track1Base;
1198 track2 = track2Base;
1199 while (track2->Neighbour(k) >= 0) {
1200 if (track1->Neighbour(k) >= 0) {
1203 track2->SetNeighbor(-1, k);
1204 track2new->SetNeighbor(-1, k ^ 1);
1206 while (track1->NextSegmentNeighbour() >= 0) {
1207 track1 = &mSectorTrackInfos[track1->NextSegmentNeighbour()];
1209 track1->SetNextSegmentNeighbour(track2new - mSectorTrackInfos);
1210 track2new->SetPrevSegmentNeighbour(track1 - mSectorTrackInfos);
1215 track1->SetNeighbor(track2->Neighbour(k), k);
1216 track2->SetNeighbor(-1, k);
1217 track2new->SetNeighbor(track1 - mSectorTrackInfos, k ^ 1);
1235 if (
Param().
par.earlyTpcTransform) {
1239 auto& cln = mConstantMem->ioPtrs.clustersNative->clustersLinear[cls.num];
1240 GPUTPCConvertImpl::convert(*mConstantMem, cls.sector, cls.row, cln.getPad(), cln.getTime(),
x,
y,
z);
1243 if (!
Param().
par.continuousTracking && CAMath::Abs(
z) > 10) {
1246 int32_t sector = track->Sector();
1247 for (int32_t attempt = 0; attempt < 2; attempt++) {
1249 const float x0 = GPUTPCGeometry::Row2X(attempt == 0 ? 63 : cls.
row);
1252 b.SetNClusters(mOutputTracks[itr].NClusters());
1253 if (CAMath::Abs(
b.Cov()[4]) >= 0.5f) {
1256 if (track->CSide()) {
1257 b.SetPar(1,
b.Par()[1] - 2 * (
z -
b.ZOffsetLinear()));
1258 b.SetZOffsetLinear(-
b.ZOffsetLinear());
1261 uint32_t
id = sector + attempt * NSECTORS;
1262 uint32_t myTrack = CAMath::AtomicAdd(&mMemory->tmpCounter[
id], 1u);
1263 mBorder[
id][myTrack] =
b;
1271 const ClusterNative* cls =
Param().par.earlyTpcTransform ? nullptr : mConstantMem->ioPtrs.clustersNative->clustersLinear;
1272 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTracks;
i += nThreads * nBlocks) {
1273 if (mOutputTracks[
i].CSide() == 0 && mTrackLinks[
i] >= 0) {
1274 if (mTrackLinks[mTrackLinks[
i]] != (int32_t)
i) {
1279 if (!trk[1]->OK() || trk[1]->CCE()) {
1282 bool celooper = (trk[0]->GetParam().GetQPt() *
Param().qptB5Scaler > 1 && trk[0]->GetParam().GetQPt() * trk[1]->GetParam().GetQPt() < 0);
1283 bool looper = trk[0]->Looper() || trk[1]->Looper() || celooper;
1284 if (!looper && trk[0]->GetParam().GetPar(3) * trk[1]->GetParam().GetPar(3) < 0) {
1288 uint32_t newRef = CAMath::AtomicAdd(&mMemory->nOutputTrackClusters, trk[0]->NClusters() + trk[1]->NClusters());
1289 if (newRef + trk[0]->NClusters() + trk[1]->NClusters() >= mNMaxOutputTrackClusters) {
1290 raiseError(GPUErrors::ERROR_MERGER_CE_HIT_OVERFLOW, newRef + trk[0]->NClusters() + trk[1]->NClusters(), mNMaxOutputTrackClusters);
1291 for (uint32_t k = newRef; k < mNMaxOutputTrackClusters; k++) {
1292 mClusters[k].num = 0;
1293 mClusters[k].state = 0;
1295 CAMath::AtomicExch(&mMemory->nOutputTrackClusters, mNMaxOutputTrackClusters);
1299 bool needswap =
false;
1302 if (
Param().
par.earlyTpcTransform) {
1303 z0max = CAMath::Max(CAMath::Abs(mClustersXYZ[trk[0]->FirstClusterRef()].
z), CAMath::Abs(mClustersXYZ[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].
z));
1304 z1max = CAMath::Max(CAMath::Abs(mClustersXYZ[trk[1]->FirstClusterRef()].
z), CAMath::Abs(mClustersXYZ[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].
z));
1306 z0max = -CAMath::Min(cls[mClusters[trk[0]->FirstClusterRef()].
num].
getTime(), cls[mClusters[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].
num].
getTime());
1307 z1max = -CAMath::Min(cls[mClusters[trk[1]->FirstClusterRef()].
num].
getTime(), cls[mClusters[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].
num].
getTime());
1309 if (z1max < z0max) {
1313 if (mClusters[trk[0]->FirstClusterRef()].
row > mClusters[trk[1]->FirstClusterRef()].
row) {
1318 GPUCommonAlgorithm::swap(trk[0], trk[1]);
1321 bool reverse[2] = {
false,
false};
1323 if (
Param().
par.earlyTpcTransform) {
1324 reverse[0] = (mClustersXYZ[trk[0]->FirstClusterRef()].z > mClustersXYZ[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].z) ^ (trk[0]->CSide() > 0);
1325 reverse[1] = (mClustersXYZ[trk[1]->FirstClusterRef()].z < mClustersXYZ[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].z) ^ (trk[1]->CSide() > 0);
1327 reverse[0] = cls[mClusters[trk[0]->FirstClusterRef()].num].getTime() < cls[mClusters[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].num].getTime();
1328 reverse[1] = cls[mClusters[trk[1]->FirstClusterRef()].num].getTime() > cls[mClusters[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].num].getTime();
1332 if (
Param().par.continuousTracking) {
1333 if (
Param().par.earlyTpcTransform) {
1334 const float z0 = trk[0]->CSide() ? CAMath::Max(mClustersXYZ[trk[0]->FirstClusterRef()].
z, mClustersXYZ[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].
z) :
CAMath::Min(mClustersXYZ[trk[0]->FirstClusterRef()].
z, mClustersXYZ[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].
z);
1335 const float z1 = trk[1]->CSide() ? CAMath::Max(mClustersXYZ[trk[1]->FirstClusterRef()].
z, mClustersXYZ[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].
z) :
CAMath::Min(mClustersXYZ[trk[1]->FirstClusterRef()].
z, mClustersXYZ[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].
z);
1336 const float offset = CAMath::Abs(z1) > CAMath::Abs(z0) ? -z0 :
z1;
1337 trk[1]->Param().Z() += trk[1]->Param().TZOffset() -
offset;
1338 trk[1]->Param().TZOffset() =
offset;
1341 const float tmax = CAMath::MaxWithRef(cls[mClusters[trk[0]->FirstClusterRef()].
num].
getTime(), cls[mClusters[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].
num].
getTime(),
1342 cls[mClusters[trk[1]->FirstClusterRef()].
num].
getTime(), cls[mClusters[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].
num].
getTime(),
1343 &mClusters[trk[0]->FirstClusterRef()], &mClusters[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1],
1344 &mClusters[trk[1]->FirstClusterRef()], &mClusters[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1], clsmax);
1345 const float offset = CAMath::Max(tmax - mConstantMem->calibObjects.fastTransformHelper->getCorrMap()->getMaxDriftTime(clsmax->
sector, clsmax->
row, cls[clsmax->
num].getPad()), 0.f);
1346 trk[1]->Param().Z() += mConstantMem->calibObjects.fastTransformHelper->getCorrMap()->convDeltaTimeToDeltaZinTimeFrame(trk[1]->CSide() * NSECTORS / 2, trk[1]->
Param().TZOffset() -
offset);
1347 trk[1]->Param().TZOffset() =
offset;
1351 int32_t
pos = newRef;
1353 int32_t lastLeg = -1;
1355 for (int32_t k = 1; k >= 0; k--) {
1356 int32_t loopstart =
reverse[k] ? (trk[k]->NClusters() - 1) : 0;
1357 int32_t loopend =
reverse[k] ? -1 : (int32_t)trk[k]->NClusters();
1358 int32_t loopinc =
reverse[k] ? -1 : 1;
1359 for (int32_t
j = loopstart;
j != loopend;
j += loopinc) {
1360 if (
Param().par.earlyTpcTransform) {
1361 mClustersXYZ[
pos] = mClustersXYZ[trk[k]->FirstClusterRef() +
j];
1363 mClusters[
pos] = mClusters[trk[k]->FirstClusterRef() +
j];
1365 if (mClusters[trk[k]->FirstClusterRef() +
j].leg != lastLeg) {
1367 lastLeg = mClusters[trk[k]->FirstClusterRef() +
j].leg;
1369 mClusters[
pos].leg =
leg;
1377 trk[1]->SetFirstClusterRef(newRef);
1378 trk[1]->SetNClusters(trk[0]->NClusters() + trk[1]->NClusters());
1383 trk[1]->SetCCE(
true);
1385 trk[1]->SetLooper(
true);
1386 trk[1]->SetLegs(
leg + 1);
1388 trk[0]->SetNClusters(0);
1389 trk[0]->SetOK(
false);
1400struct GPUTPCGMMerger_CompareClusterIdsLooper {
1401 struct clcomparestruct {
1410 GPUd() bool operator()(const int16_t
aa, const int16_t
bb)
1412 const clcomparestruct&
a =
cmp2[
aa];
1416 if (
a.leg !=
b.leg) {
1417 return ((
leg > 0) ^ (
a.leg >
b.leg));
1426struct GPUTPCGMMerger_CompareClusterIds {
1429 GPUd() bool operator()(const int16_t
aa, const int16_t
bb)
1433 if (
a.row !=
b.row) {
1434 return (
a.row >
b.row);
1442GPUd()
void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1446 for (int32_t itr = iBlock * nThreads + iThread; itr < SectorTrackInfoLocalTotal(); itr += nThreads * nBlocks) {
1450 if (track.PrevSegmentNeighbour() >= 0) {
1453 if (track.PrevNeighbour() >= 0) {
1460 tr->SetPrevSegmentNeighbour(1000000000);
1462 if (nParts >= kMaxParts) {
1465 if (nHits + tr->NClusters() > kMaxClusters) {
1468 nHits += tr->NClusters();
1471 trackParts[nParts++] = tr;
1472 for (int32_t
i = 0;
i < 2;
i++) {
1473 if (tr->ExtrapolatedTrackId(
i) != -1) {
1474 if (nParts >= kMaxParts) {
1477 if (nHits + mSectorTrackInfos[tr->ExtrapolatedTrackId(
i)].NClusters() > kMaxClusters) {
1480 trackParts[nParts] = &mSectorTrackInfos[tr->ExtrapolatedTrackId(
i)];
1481 trackParts[nParts++]->SetLeg(
leg);
1482 nHits += mSectorTrackInfos[tr->ExtrapolatedTrackId(
i)].NClusters();
1485 int32_t jtr = tr->NextSegmentNeighbour();
1487 tr = &(mSectorTrackInfos[jtr]);
1488 tr->SetPrevSegmentNeighbour(1000000002);
1491 jtr = trbase->NextNeighbour();
1493 trbase = &(mSectorTrackInfos[jtr]);
1495 if (tr->PrevSegmentNeighbour() >= 0) {
1498 tr->SetPrevSegmentNeighbour(1000000001);
1506 if (nParts > 1 &&
leg == 0) {
1509 if (
a->X() !=
b->X()) {
1510 return (a->X() > b->X());
1512 if (
a->Y() !=
b->Y()) {
1513 return (a->Y() > b->Y());
1515 if (
a->Z() !=
b->Z()) {
1516 return (a->Z() > b->Z());
1518 return a->QPt() >
b->QPt();
1520 return (
a->X() >
b->X());
1525 if (
Param().rec.tpc.dropLoopers &&
leg > 0) {
1530 trackCluster trackClusters[kMaxClusters];
1532 for (int32_t ipart = 0; ipart < nParts; ipart++) {
1534 CADEBUG(printf(
"Collect Track %d Part %d QPt %f DzDs %f\n", mMemory->nOutputTracks, ipart, t->QPt(), t->DzDs()));
1535 int32_t nTrackHits = t->NClusters();
1536 trackCluster*
c2 = trackClusters + nHits + nTrackHits - 1;
1537 for (int32_t
i = 0;
i < nTrackHits;
i++,
c2--) {
1538 const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()];
1539 const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() +
i];
1540 uint32_t
id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0];
1541 *
c2 = trackCluster{
id, (
uint8_t)ic.RowIndex(), t->Sector(), t->Leg()};
1543 nHits += nTrackHits;
1551 for (int32_t
i = 1;
i < nHits;
i++) {
1552 if (trackClusters[
i].
row > trackClusters[
i - 1].
row || trackClusters[
i].
id == trackClusters[
i - 1].
id) {
1558 int32_t firstTrackIndex = 0;
1559 int32_t lastTrackIndex = nParts - 1;
1561 int32_t nTmpHits = 0;
1562 trackCluster trackClustersUnsorted[kMaxClusters];
1563 int16_t clusterIndices[kMaxClusters];
1564 for (int32_t
i = 0;
i < nHits;
i++) {
1565 trackClustersUnsorted[
i] = trackClusters[
i];
1566 clusterIndices[
i] =
i;
1573 for (int32_t
i = 0;
i < nParts;
i++) {
1574 if (trackParts[
i]->Leg() == 0 || trackParts[
i]->Leg() ==
leg) {
1576 if (
Param().par.earlyTpcTransform) {
1577 zt = CAMath::Min(CAMath::Abs(trackParts[
i]->ClusterZT0()), CAMath::Abs(trackParts[
i]->ClusterZTN()));
1579 zt = -trackParts[
i]->MinClusterZT();
1583 baseLeg = trackParts[
i]->Leg();
1587 int32_t iLongest = 1e9;
1589 for (int32_t
i = (baseLeg ? (nParts - 1) : 0); baseLeg ? (
i >= 0) : (
i < nParts); baseLeg ?
i-- :
i++) {
1590 if (trackParts[
i]->Leg() != baseLeg) {
1593 if (trackParts[
i]->OrigTrack()->NHits() >
length) {
1595 length = trackParts[
i]->OrigTrack()->NHits();
1599 if (
Param().par.earlyTpcTransform) {
1600 outwards = (trackParts[iLongest]->ClusterZT0() > trackParts[iLongest]->ClusterZTN()) ^ trackParts[iLongest]->CSide();
1602 outwards = trackParts[iLongest]->ClusterZT0() < trackParts[iLongest]->ClusterZTN();
1604 GPUTPCGMMerger_CompareClusterIdsLooper::clcomparestruct clusterSort[kMaxClusters];
1605 for (int32_t iPart = 0; iPart < nParts; iPart++) {
1607 int32_t nTrackHits = t->NClusters();
1608 for (int32_t
j = 0;
j < nTrackHits;
j++) {
1609 int32_t
i = nTmpHits +
j;
1610 clusterSort[
i].leg = t->Leg();
1612 nTmpHits += nTrackHits;
1615 GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIdsLooper(baseLeg,
outwards, trackClusters, clusterSort));
1617 GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIds(trackClusters));
1620 firstTrackIndex = lastTrackIndex = -1;
1621 for (int32_t
i = 0;
i < nParts;
i++) {
1622 nTmpHits += trackParts[
i]->NClusters();
1623 if (nTmpHits > clusterIndices[0] && firstTrackIndex == -1) {
1624 firstTrackIndex =
i;
1626 if (nTmpHits > clusterIndices[nHits - 1] && lastTrackIndex == -1) {
1631 int32_t nFilteredHits = 0;
1632 int32_t indPrev = -1;
1633 for (int32_t
i = 0;
i < nHits;
i++) {
1634 int32_t ind = clusterIndices[
i];
1635 if (indPrev >= 0 && trackClustersUnsorted[ind].
id == trackClustersUnsorted[indPrev].
id) {
1639 trackClusters[nFilteredHits] = trackClustersUnsorted[ind];
1642 nHits = nFilteredHits;
1645 const uint32_t iOutTrackFirstCluster = CAMath::AtomicAdd(&mMemory->nOutputTrackClusters, (uint32_t)nHits);
1646 if (iOutTrackFirstCluster >= mNMaxOutputTrackClusters) {
1647 raiseError(GPUErrors::ERROR_MERGER_HIT_OVERFLOW, iOutTrackFirstCluster, mNMaxOutputTrackClusters);
1648 CAMath::AtomicExch(&mMemory->nOutputTrackClusters, mNMaxOutputTrackClusters);
1654 for (int32_t
i = 0;
i < nHits;
i++) {
1656 if (
Param().
par.earlyTpcTransform) {
1657 const GPUTPCClusterData&
c = GetConstantMem()->tpcTrackers[trackClusters[
i].sector].ClusterData()[trackClusters[
i].id - GetConstantMem()->tpcTrackers[trackClusters[
i].sector].Data().ClusterIdOffset()];
1662 clXYZ[
i].
amp =
c.amp;
1663#ifdef GPUCA_TPC_RAW_PROPAGATE_PAD_ROW_TIME
1664 clXYZ[
i].pad =
c.mPad;
1665 clXYZ[
i].time =
c.mTime;
1669 const ClusterNative&
c = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[trackClusters[
i].id];
1673 cl[
i].
row = trackClusters[
i].row;
1674 cl[
i].
num = trackClusters[
i].id;
1675 cl[
i].
sector = trackClusters[
i].sector;
1676 cl[
i].
leg = trackClusters[
i].leg;
1679 uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nOutputTracks, 1u);
1680 if (iOutputTrack >= mNMaxTracks) {
1681 raiseError(GPUErrors::ERROR_MERGER_TRACK_OVERFLOW, iOutputTrack, mNMaxTracks);
1682 CAMath::AtomicExch(&mMemory->nOutputTracks, mNMaxTracks);
1688 mergedTrack.SetFlags(0);
1689 mergedTrack.SetOK(1);
1690 mergedTrack.SetLooper(
leg > 0);
1691 mergedTrack.SetLegs(
leg);
1692 mergedTrack.SetNClusters(nHits);
1693 mergedTrack.SetFirstClusterRef(iOutTrackFirstCluster);
1696 mergedTrack.SetCSide(
p2.CSide());
1699 const float toX =
Param().par.earlyTpcTransform ? mClustersXYZ[iOutTrackFirstCluster].x : GPUTPCGeometry::Row2X(cl[0].
row);
1702 p1.Y() =
b.Par()[0];
1703 p1.Z() =
b.Par()[1];
1704 p1.SinPhi() =
b.Par()[2];
1709 p1.SinPhi() =
p2.SinPhi();
1711 p1.TZOffset() =
p2.TZOffset();
1712 p1.DzDs() =
p2.DzDs();
1713 p1.QPt() =
p2.QPt();
1714 mergedTrack.SetAlpha(
p2.Alpha());
1715 if (CAMath::Abs(
Param().polynomialField.GetNominalBz()) < (gpu_common_constants::kZeroFieldCut * gpu_common_constants::kCLight)) {
1716 p1.QPt() = 100.f /
Param().rec.bz0Pt10MeV;
1729 if (
Param().
par.earlyTpcTransform) {
1731 CEside = (mergedTrack.CSide() != 0) ^ (clXYZ[0].
z > clXYZ[nHits - 1].
z);
1733 auto& cls = mConstantMem->ioPtrs.clustersNative->clustersLinear;
1734 CEside = cls[cl[0].
num].getTime() < cls[cl[nHits - 1].
num].getTime();
1736 MergeCEFill(trackParts[CEside ? lastTrackIndex : firstTrackIndex], cl[CEside ? (nHits - 1) : 0],
Param().
par.earlyTpcTransform ? &(mClustersXYZ + iOutTrackFirstCluster)[CEside ? (nHits - 1) : 0] : nullptr, iOutputTrack);
1741GPUd()
void GPUTPCGMMerger::SortTracksPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1743 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTracks;
i += nThreads * nBlocks) {
1744 mTrackOrderProcess[
i] =
i;
1748GPUd()
void GPUTPCGMMerger::PrepareClustersForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1750 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTracks;
i += nBlocks * nThreads) {
1755GPUd()
void GPUTPCGMMerger::SortTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1757#ifndef GPUCA_SPECIALIZE_THRUST_SORTS
1758 if (iThread || iBlock) {
1762 auto comp = [
cmp = mOutputTracks](
const int32_t
aa,
const int32_t
bb) {
1765 if (
a.CCE() !=
b.CCE()) {
1766 return a.CCE() >
b.CCE();
1768 if (
a.Legs() !=
b.Legs()) {
1769 return a.Legs() >
b.Legs();
1772 if (
a.NClusters() !=
b.NClusters()) {
1773 return a.NClusters() > b.NClusters();
1774 }
if (CAMath::Abs(
a.GetParam().GetQPt()) != CAMath::Abs(
b.GetParam().GetQPt())) {
1775 return CAMath::Abs(a.GetParam().GetQPt()) > CAMath::Abs(b.GetParam().GetQPt());
1776 }
if (
a.GetParam().GetY() !=
b.GetParam().GetY()) {
1777 return a.GetParam().GetY() > b.GetParam().GetY();
1781 return a.NClusters() >
b.NClusters();
1785 GPUCommonAlgorithm::sortDeviceDynamic(mTrackOrderProcess, mTrackOrderProcess + mMemory->nOutputTracks, comp);
1789GPUd()
void GPUTPCGMMerger::SortTracksQPt(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1791#ifndef GPUCA_SPECIALIZE_THRUST_SORTS
1792 if (iThread || iBlock) {
1796 auto comp = [
cmp = mOutputTracks](
const int32_t
aa,
const int32_t
bb) {
1801 return CAMath::Abs(
a.GetParam().GetQPt()) > CAMath::Abs(
b.GetParam().GetQPt());
1802 }
if (
a.GetParam().GetY() !=
b.GetParam().GetY()) {
1803 return a.GetParam().GetY() >
b.GetParam().GetY();
1805 return a.GetParam().GetZ() >
b.GetParam().GetZ();
1807 return CAMath::Abs(
a.GetParam().GetQPt()) > CAMath::Abs(
b.GetParam().GetQPt());
1811 GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nOutputTracks, comp);
1815GPUd()
void GPUTPCGMMerger::PrepareClustersForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1817 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTracks;
i += nBlocks * nThreads) {
1818 mTrackOrderAttach[mTrackSort[
i]] =
i;
1821 for (uint32_t
j = 0;
j < trk.NClusters();
j++) {
1823 CAMath::AtomicAdd(&mSharedCount[mClusters[trk.FirstClusterRef() +
j].num], 1u);
1829GPUd()
void GPUTPCGMMerger::PrepareClustersForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1831 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTrackClusters;
i += nBlocks * nThreads) {
1832 if (mSharedCount[mClusters[
i].
num] > 1) {
1836 if (mClusterStateExt) {
1837 for (uint32_t
i = iBlock * nThreads + iThread;
i < mNMaxClusters;
i += nBlocks * nThreads) {
1838 uint8_t state = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[
i].getFlags();
1839 if (mSharedCount[
i] > 1) {
1842 mClusterStateExt[
i] =
state;
1849 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTracks;
i += nThreads * nBlocks) {
1850 mTrackSort[mTrackOrderAttach[
i]] =
i;
1852 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTrackClusters;
i += nThreads * nBlocks) {
1853 mClusterAttachment[mClusters[
i].num] = 0;
1859 for (uint32_t
i = iBlock * nThreads + iThread;
i < mMemory->nOutputTracks;
i += nThreads * nBlocks) {
1861 if (!trk.OK() || trk.NClusters() == 0) {
1864 uint8_t goodLeg = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].leg;
1865 for (uint32_t
j = 0;
j < trk.NClusters();
j++) {
1866 int32_t
id = mClusters[trk.FirstClusterRef() +
j].num;
1868 uint8_t clusterState = mClusters[trk.FirstClusterRef() +
j].state;
1874 if (mClusters[trk.FirstClusterRef() +
j].leg == goodLeg) {
1877 CAMath::AtomicMax(&mClusterAttachment[
id],
weight);
1884 for (uint32_t
i = iBlock * nThreads + iThread;
i < mNMaxClusters;
i += nThreads * nBlocks) {
1885 if (mClusterAttachment[
i] != 0) {
1891GPUd()
void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1893 const float lowPtThresh =
Param().rec.tpc.rejectQPtB5 * 1.1f;
1895 const auto& trk = mOutputTracks[
i];
1896 const auto& p = trk.GetParam();
1897 const float qptabs = CAMath::Abs(p.GetQPt());
1898 if (trk.NClusters() && qptabs *
Param().qptB5Scaler > 5.f && qptabs *
Param().qptB5Scaler <= lowPtThresh) {
1899 const int32_t sector = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].sector;
1900 const float refz = p.GetZ() + (
Param().par.earlyTpcTransform ? p.GetTZOffset() : GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convVertexTimeToZOffset(sector, p.GetTZOffset(),
Param().continuousMaxTimeBin)) + (trk.CSide() ? -100 : 100);
1902 CAMath::SinCos(trk.GetAlpha(), sinA, cosA);
1903 float gx = cosA * p.GetX() - sinA * p.GetY();
1904 float gy = cosA * p.GetY() + sinA * p.GetX();
1905 float bz =
Param().polynomialField.GetFieldBz(gx, gy, p.GetZ());
1906 const float r1 = p.GetQPt() * bz;
1907 const float r = CAMath::Abs(r1) > 0.0001f ? (1.f / r1) : 10000;
1908 const float mx = p.GetX() +
r * p.GetSinPhi();
1909 const float my = p.GetY() -
r * CAMath::Sqrt(1 - p.GetSinPhi() * p.GetSinPhi());
1910 const float gmx = cosA * mx - sinA * my;
1911 const float gmy = cosA * my + sinA * mx;
1912 uint32_t myId = CAMath::AtomicAdd(&mMemory->nLooperMatchCandidates, 1u);
1913 if (myId >= mNMaxLooperMatches) {
1914 raiseError(GPUErrors::ERROR_LOOPER_MATCH_OVERFLOW, myId, mNMaxLooperMatches);
1915 CAMath::AtomicExch(&mMemory->nLooperMatchCandidates, mNMaxLooperMatches);
1943GPUd()
void GPUTPCGMMerger::MergeLoopersSort(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1945#ifndef GPUCA_SPECIALIZE_THRUST_SORTS
1946 if (iThread || iBlock) {
1950 GPUCommonAlgorithm::sortDeviceDynamic(mLooperCandidates, mLooperCandidates + mMemory->nLooperMatchCandidates, comp);
1954GPUd()
void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1958#if GPUCA_MERGE_LOOPER_MC && !defined(GPUCA_GPUCODE)
1959 std::vector<int64_t> paramLabels(mMemory->nLooperMatchCandidates);
1960 for (uint32_t
i = 0;
i < mMemory->nLooperMatchCandidates;
i++) {
1961 paramLabels[
i] = GetTrackLabel(mOutputTracks[
params[
i].
id]);
1967 if (!mRec->GetProcessingSettings().runQA) {
1968 throw std::runtime_error(
"Need QA enabled for the Merge Loopers MC QA");
1973 for (uint32_t
j =
i + 1;
j < mMemory->nLooperMatchCandidates;
j++) {
1975 if (CAMath::Abs(
params[
j].refz) > CAMath::Abs(
params[
i].refz) + 100.f) {
1983 const auto& trk1 = mOutputTracks[
params[
i].id];
1984 const auto& trk2 = mOutputTracks[
params[
j].id];
1985 const auto& param1 = trk1.GetParam();
1986 const auto& param2 = trk2.GetParam();
1987 if (CAMath::Abs(param1.GetDzDs()) > 0.03f && CAMath::Abs(param2.GetDzDs()) > 0.03f && param1.GetDzDs() * param2.GetDzDs() * param1.GetQPt() * param2.GetQPt() < 0) {
1992 const float dznormalized = (CAMath::Abs(
params[
j].refz) - CAMath::Abs(
params[
i].refz)) / (CAMath::TwoPi() * 0.5f * (CAMath::Abs(param1.GetDzDs()) + CAMath::Abs(param2.GetDzDs())) * 1.f / (0.5f * (CAMath::Abs(param1.GetQPt()) + CAMath::Abs(param2.GetQPt())) * CAMath::Abs(
Param().polynomialField.GetNominalBz())));
1993 const float phasecorr = CAMath::Modf((CAMath::ASin(param1.GetSinPhi()) + trk1.GetAlpha() - CAMath::ASin(param2.GetSinPhi()) - trk2.GetAlpha()) / CAMath::TwoPi() + 5.5f, 1.f) - 0.5f;
1994 const float phasecorrdirection = (
params[
j].refz * param1.GetQPt() * param1.GetDzDs()) > 0 ? 1 : -1;
1995 const float dzcorr = dznormalized + phasecorr * phasecorrdirection;
1996 const bool sameside = !(trk1.CSide() ^ trk2.CSide());
1997 const float dzcorrlimit[4] = {sameside ? 0.018f : 0.012f, sameside ? 0.12f : 0.025f, 0.14f, 0.15f};
1998 const int32_t dzcorrcount = sameside ? 4 : 2;
1999 bool dzcorrok =
false;
2001 for (int32_t k = 0; k < dzcorrcount; k++) {
2002 const float d = CAMath::Abs(dzcorr - 0.5f * k);
2003 if (d <= dzcorrlimit[k]) {
2005 dznorm = d / dzcorrlimit[k];
2014 const float dtgl = param1.GetDzDs() - (param1.GetQPt() * param2.GetQPt() > 0 ? param2.GetDzDs() : -param2.GetDzDs());
2015 const float dqpt = (CAMath::Abs(param1.GetQPt()) - CAMath::Abs(param2.GetQPt())) / CAMath::Min(param1.GetQPt(), param2.GetQPt());
2016 float d = CAMath::Sum2(dtgl * (1.f / 0.03f), dqpt * (1.f / 0.04f)) + d2xy * (1.f / 4.f) + dznorm * (1.f / 0.3f);
2018#if GPUCA_MERGE_LOOPER_MC && !defined(GPUCA_GPUCODE)
2019 const int64_t label1 = paramLabels[
i];
2020 const int64_t label2 = paramLabels[
j];
2021 bool labelEQ = label1 != -1 && label1 == label2;
2022 if (1 || EQ || labelEQ) {
2024 static auto& tup =
GPUROOTDump<TNtuple>::get(
"mergeloopers",
"labeleq:sides:d2xy:tgl1:tgl2:qpt1:qpt2:dz:dzcorr:dtgl:dqpt:dznorm:bs");
2025 tup.Fill((
float)labelEQ, (trk1.CSide() ? 1 : 0) | (trk2.CSide() ? 2 : 0), d2xy, param1.GetDzDs(), param2.GetDzDs(), param1.GetQPt(), param2.GetQPt(),
CAMath::
Abs(
params[
j].refz) -
CAMath::
Abs(
params[
i].refz), dzcorr, dtgl, dqpt, dznorm, bs);
2026 static auto tup2 =
GPUROOTDump<TNtuple>::getNew(
"mergeloopers2",
"labeleq:refz1:refz2:tgl1:tgl2:qpt1:qpt2:snp1:snp2:a1:a2:dzn:phasecor:phasedir:dzcorr");
2027 tup2.Fill((
float)labelEQ,
params[
i].refz,
params[
j].refz, param1.GetDzDs(), param2.GetDzDs(), param1.GetQPt(), param2.GetQPt(), param1.GetSinPhi(), param2.GetSinPhi(), trk1.GetAlpha(), trk2.GetAlpha(), dznormalized, phasecorr, phasecorrdirection, dzcorr);
2041 mOutputTracks[
params[
j].id].SetMergedLooper(
true);
2042 if (CAMath::Abs(param2.GetQPt() *
Param().qptB5Scaler) >=
Param().
rec.tpc.rejectQPtB5) {
2043 mOutputTracks[
params[
i].id].SetMergedLooper(
true);
Class of a TPC cluster in TPC-native coordinates (row, time)
A const (ready only) version of MCTruthContainer.
#define get_global_size(dim)
#define get_global_id(dim)
#define GPUCA_DEBUG_STREAMER_CHECK(...)
#define GPUCA_DETERMINISTIC_CODE(det, indet)
#define GPUCA_MERGER_MAX_TRACK_CLUSTERS
#define GPUCA_MAX_SIN_PHI_LOW
#define GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(QPTB5)
#define GPUCA_MAX_SIN_PHI
#define CADEBUG2(cmd,...)
const GPUTPCGMMerger::trackCluster & b1
const GPUTPCGMMerger::trackCluster & a1
const GPUTPCGMMerger::trackCluster *const mCmp
const clcomparestruct *const cmp2
const GPUTPCGMMerger::trackCluster *const cmp1
bool const GPUTPCGMMerger::trackCluster * c1
bool const GPUTPCGMMerger::trackCluster const clcomparestruct * c2
constexpr int p1()
constexpr to accelerate the coordinates changing
std::array< int, 64 > reverse(std::array< int, 64 > a)
Class for time synchronization of RawReader instances.
static void computePointerWithAlignment(T *&basePtr, S *&objPtr, size_t nEntries=1)
void AllocateAndInitializeLate()
static GPUROOTDump< T, Args... > getNew(const char *name1, Names... names)
static GPUROOTDump< T, Args... > & get(const char *name1, Names... names)
RecoStepField GetRecoStepsGPU() const
int16_t RegisterMemoryAllocation(T *proc, void *(T::*setPtr)(void *), int32_t type, const char *name="", const GPUMemoryReuse &re=GPUMemoryReuse())
RecoStepField GetRecoSteps() const
const GPUParam & GetParam() const
const GPUConstantMem & GetConstantMem() const
GPUMemorySizeScalers * MemoryScalers()
const GPUSettingsProcessing & GetProcessingSettings() const
virtual const GPUDefParameters & getGPUParameters(bool doGPU) const =0
void * SetPointersRefitScratch(void *mem)
void RegisterMemoryAllocation()
void * SetPointersMemory(void *mem)
static constexpr const int32_t NSECTORS
void * SetPointersOutput(void *mem)
void InitializeProcessor()
void * SetPointersOutputO2MC(void *mem)
void * SetPointersOutputO2Scratch(void *mem)
void * SetPointersMerger(void *mem)
void * SetPointersOutputO2Clus(void *mem)
void * SetPointersOutputState(void *mem)
void SetMaxData(const GPUTrackingInOutPointers &io)
void * SetPointersOutputO2(void *mem)
int32_t int32_t int32_t bool float float Z
int32_t int32_t int32_t bool float Y
GLint GLint GLsizei GLint border
GLfloat GLfloat GLfloat alpha
GLuint GLuint GLfloat weight
GLboolean GLboolean GLboolean b
GLenum const GLfloat * params
GLuint GLsizei GLsizei * length
typedef void(APIENTRYP PFNGLCULLFACEPROC)(GLenum mode)
GLboolean GLboolean GLboolean GLboolean a
GLdouble GLdouble GLdouble z
uint8_t itsSharedClusterMap uint8_t
Node par(int index)
Parameters.
GPUd() const expr uint32_t MultivariatePolynomialHelper< Dim
float float float float z1
if(!okForPhiMin(phi0, phi1))
Global TPC definitions and constants.
@ streamMergeBorderTracksBest
stream MergeBorderTracks best track
@ streamMergeBorderTracksAll
stream MergeBorderTracks all tracks
std::string getTime(uint64_t ts)
GPUTPCTracker tpcTrackers[GPUCA_NSECTORS]
size_t getValue(size_t maxVal, size_t val)
size_t NTPCMergedTrackHits(size_t tpcSectorTrackHitss)
size_t NTPCMergedTracks(size_t tpcSectorTracks)
@ clustererAndSharedFlags
const o2::tpc::ClusterNativeAccess * clustersNative
unsigned int nClustersTotal
char const *restrict const cmp