Skip to content

Commit 6c9dc79

Browse files
davidrohrdavidrohr
committed
GPU TPC: Shift all segments of looping tracks once before track fit
1 parent caa743a commit 6c9dc79

File tree

11 files changed

+160
-144
lines changed

11 files changed

+160
-144
lines changed

‎GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h‎

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -57,9 +57,9 @@
5757
#define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
5858
#define GPUCA_LB_GPUTPCGMMergerCollect 512
5959
#define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
60-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
61-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
62-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
60+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
61+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
62+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
6363
#define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
6464
#define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
6565
#define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -119,9 +119,9 @@
119119
#define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
120120
#define GPUCA_LB_GPUTPCGMMergerCollect 512
121121
#define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
122-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
123-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
124-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
122+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
123+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
124+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
125125
#define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
126126
#define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
127127
#define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -181,9 +181,9 @@
181181
#define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
182182
#define GPUCA_LB_GPUTPCGMMergerCollect 256, 2
183183
#define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
184-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
185-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
186-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
184+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
185+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
186+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
187187
#define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
188188
#define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
189189
#define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -243,9 +243,9 @@
243243
#define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256
244244
#define GPUCA_LB_GPUTPCGMMergerCollect 128, 2
245245
#define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
246-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
247-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
248-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
246+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
247+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
248+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
249249
#define GPUCA_LB_GPUTPCGMMergerFinalize_0 256
250250
#define GPUCA_LB_GPUTPCGMMergerFinalize_1 256
251251
#define GPUCA_LB_GPUTPCGMMergerFinalize_2 256
@@ -398,14 +398,14 @@
398398
#ifndef GPUCA_LB_GPUTPCGMMergerSortTracksPrepare
399399
#define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256
400400
#endif
401-
#ifndef GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0
402-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256
401+
#ifndef GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0
402+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step0 256
403403
#endif
404-
#ifndef GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1
405-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step1 256
404+
#ifndef GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1
405+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step1 256
406406
#endif
407-
#ifndef GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2
408-
#define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step2 256
407+
#ifndef GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2
408+
#define GPUCA_LB_GPUTPCGMMergerPrepareForFit_step2 256
409409
#endif
410410
#ifndef GPUCA_LB_GPUTPCGMMergerFinalize_step0
411411
#define GPUCA_LB_GPUTPCGMMergerFinalize_step0 256

‎GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -196,11 +196,11 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
196196
}
197197
runKernel<GPUMemClean16>({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.SharedCount(), maxId * sizeof(*MergerShadowAll.SharedCount()));
198198
runKernel<GPUMemClean16>({{numBlocks, -ThreadCount(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.ClusterAttachment(), maxId * sizeof(*MergerShadowAll.ClusterAttachment()));
199-
runKernel<GPUTPCGMMergerPrepareClusters, 0>(GetGridAuto(0, deviceType));
199+
runKernel<GPUTPCGMMergerPrepareForFit, 0>(GetGridAuto(0, deviceType));
200200
CondWaitEvent(waitForTransfer, &mEvents->single);
201201
runKernel<GPUTPCGMMergerSortTracksQPt>(GetGridAuto(0, deviceType));
202-
runKernel<GPUTPCGMMergerPrepareClusters, 1>(GetGridAuto(0, deviceType));
203-
runKernel<GPUTPCGMMergerPrepareClusters, 2>(GetGridAuto(0, deviceType));
202+
runKernel<GPUTPCGMMergerPrepareForFit, 1>(GetGridAuto(0, deviceType));
203+
runKernel<GPUTPCGMMergerPrepareForFit, 2>(GetGridAuto(0, deviceType));
204204

205205
DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingPrepareFit, doGPU, Merger, &GPUTPCGMMerger::DumpFitPrepare, *mDebugFile);
206206

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 24 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -325,7 +325,7 @@ void* GPUTPCGMMerger::SetPointersMerger(void* mem)
325325
computePointerWithAlignment(mem, mTrackIDs, GPUCA_NSECTORS * mNMaxSingleSectorTracks); // UnpackResetIds - RefitSectorTracks - UnpackSectorGlobal
326326
memMax = (void*)std::max((size_t)mem, (size_t)memMax);
327327
mem = memBase;
328-
computePointerWithAlignment(mem, mTrackSort, mNMaxTracks); // PrepareClustersForFit0 - SortTracksQPt - PrepareClustersForFit1 - PrepareClustersForFit1 / Finalize0 - Finalize2
328+
computePointerWithAlignment(mem, mTrackSort, mNMaxTracks); // PrepareForFit0 - SortTracksQPt - PrepareForFit1 - PrepareForFit1 / Finalize0 - Finalize2
329329
computePointerWithAlignment(mem, mSharedCount, mNMaxClusters);
330330
memMax = (void*)std::max((size_t)mem, (size_t)memMax);
331331
mem = memBase;
@@ -492,7 +492,8 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
492492
trk.DzDs() = inTrack->Param().GetDzDs();
493493
trk.QPt() = inTrack->Param().GetQPt();
494494
trk.TZOffset() = GetConstantMem()->calibObjects.fastTransformHelper->getCorrMap()->convZOffsetToVertexTime(sector, inTrack->Param().GetZOffset(), Param().continuousMaxTimeBin);
495-
trk.ShiftZ(this, sector, sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN(), inTrack->Param().GetX(), inTrack->Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
495+
const auto tmp = sectorTrack.ClusterZTN() > sectorTrack.ClusterZT0() ? std::array<float, 2>{sectorTrack.ClusterZTN(), sectorTrack.ClusterZT0()} : std::array<float, 2>{sectorTrack.ClusterZT0(), sectorTrack.ClusterZTN()};
496+
trk.ShiftZ(this, sector, tmp[0], tmp[1], inTrack->Param().GetX()); // We do not store the inner / outer cluster X, so we just use the track X instead
496497
sectorTrack.SetX2(0.f);
497498
for (int32_t way = 0; way < 2; way++) {
498499
if (way) {
@@ -514,7 +515,7 @@ GPUd() int32_t GPUTPCGMMerger::RefitSectorTrack(GPUTPCGMSectorTrack& sectorTrack
514515
const ClusterNative& cl = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[GetConstantMem()->ioPtrs.clustersNative->clusterOffset[sector][0] + clusterIndex];
515516
flags = cl.getFlags();
516517
GetConstantMem()->calibObjects.fastTransformHelper->Transform(sector, row, cl.getPad(), cl.getTime(), x, y, z, trk.TZOffset());
517-
if (prop.PropagateToXAlpha(x, alpha, true)) {
518+
if (prop.PropagateToXAlpha(x, alpha, way == 0)) {
518519
return way == 0;
519520
}
520521
trk.ConstrainSinPhi();
@@ -1738,7 +1739,7 @@ GPUd() void GPUTPCGMMerger::SortTracksPrepare(int32_t nBlocks, int32_t nThreads,
17381739
}
17391740
}
17401741

1741-
GPUd() void GPUTPCGMMerger::PrepareClustersForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1742+
GPUd() void GPUTPCGMMerger::PrepareForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17421743
{
17431744
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
17441745
mTrackSort[i] = i;
@@ -1763,21 +1764,37 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int32_t nBlocks, int32_t nThreads, int
17631764
#endif
17641765
}
17651766

1766-
GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1767+
GPUd() void GPUTPCGMMerger::PrepareForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17671768
{
17681769
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
17691770
mTrackOrderAttach[mTrackSort[i]] = i;
1770-
const GPUTPCGMMergedTrack& trk = mMergedTracks[i];
1771+
GPUTPCGMMergedTrack& trk = mMergedTracks[i];
17711772
if (trk.OK()) {
17721773
for (uint32_t j = 0; j < trk.NClusters(); j++) {
17731774
mClusterAttachment[mClusters[trk.FirstClusterRef() + j].num] = attachAttached | attachGood;
17741775
CAMath::AtomicAdd(&mSharedCount[mClusters[trk.FirstClusterRef() + j].num], 1u);
17751776
}
1777+
if (!trk.CCE() && !trk.MergedLooper()) {
1778+
GPUTPCGMMergedTrack* updTrk = &trk;
1779+
while (updTrk->PrevSegment() >= 0) {
1780+
updTrk = &mMergedTracks[updTrk->PrevSegment()];
1781+
}
1782+
const auto &cl0 = mClusters[trk.FirstClusterRef()], &cln = mClusters[updTrk->FirstClusterRef() + updTrk->NClusters() - 1];
1783+
const auto& GPUrestrict() cls = GetConstantMem()->ioPtrs.clustersNative->clustersLinear;
1784+
float z0 = cls[cl0.num].getTime(), zn = cls[cln.num].getTime();
1785+
const auto tmp = zn > z0 ? std::array<float, 3>{zn, z0, GPUTPCGeometry::Row2X(cln.row)} : std::array<float, 3>{z0, zn, GPUTPCGeometry::Row2X(cl0.row)};
1786+
trk.Param().ShiftZ(this, cl0.sector, tmp[0], tmp[1], tmp[2]);
1787+
updTrk = &trk;
1788+
while (updTrk->PrevSegment() >= 0) {
1789+
updTrk = &mMergedTracks[updTrk->PrevSegment()];
1790+
updTrk->Param().TZOffset() = trk.Param().TZOffset();
1791+
}
1792+
}
17761793
}
17771794
}
17781795
}
17791796

1780-
GPUd() void GPUTPCGMMerger::PrepareClustersForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1797+
GPUd() void GPUTPCGMMerger::PrepareForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17811798
{
17821799
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTrackClusters; i += nBlocks * nThreads) {
17831800
if (mSharedCount[mClusters[i].num] > 1) {

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -170,9 +170,9 @@ class GPUTPCGMMerger : public GPUProcessor
170170
GPUd() void SortTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
171171
GPUd() void SortTracksQPt(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
172172
GPUd() void SortTracksPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
173-
GPUd() void PrepareClustersForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
174-
GPUd() void PrepareClustersForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
175-
GPUd() void PrepareClustersForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
173+
GPUd() void PrepareForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
174+
GPUd() void PrepareForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
175+
GPUd() void PrepareForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
176176
GPUd() void LinkExtrapolatedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
177177
GPUd() void CollectMergedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);
178178
GPUd() void Finalize0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread);

‎GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -163,21 +163,21 @@ GPUdii() void GPUTPCGMMergerSortTracksPrepare::Thread<0>(int32_t nBlocks, int32_
163163
}
164164

165165
template <>
166-
GPUdii() void GPUTPCGMMergerPrepareClusters::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
166+
GPUdii() void GPUTPCGMMergerPrepareForFit::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
167167
{
168-
merger.PrepareClustersForFit0(nBlocks, nThreads, iBlock, iThread);
168+
merger.PrepareForFit0(nBlocks, nThreads, iBlock, iThread);
169169
}
170170

171171
template <>
172-
GPUdii() void GPUTPCGMMergerPrepareClusters::Thread<1>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
172+
GPUdii() void GPUTPCGMMergerPrepareForFit::Thread<1>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
173173
{
174-
merger.PrepareClustersForFit1(nBlocks, nThreads, iBlock, iThread);
174+
merger.PrepareForFit1(nBlocks, nThreads, iBlock, iThread);
175175
}
176176

177177
template <>
178-
GPUdii() void GPUTPCGMMergerPrepareClusters::Thread<2>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
178+
GPUdii() void GPUTPCGMMergerPrepareForFit::Thread<2>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger)
179179
{
180-
merger.PrepareClustersForFit2(nBlocks, nThreads, iBlock, iThread);
180+
merger.PrepareForFit2(nBlocks, nThreads, iBlock, iThread);
181181
}
182182

183183
template <>

‎GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -138,7 +138,7 @@ class GPUTPCGMMergerCollect : public GPUTPCGMMergerGeneral
138138
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& merger);
139139
};
140140

141-
class GPUTPCGMMergerPrepareClusters : public GPUTPCGMMergerGeneral
141+
class GPUTPCGMMergerPrepareForFit : public GPUTPCGMMergerGeneral
142142
{
143143
public:
144144
template <int32_t iKernel = defaultKernel>

‎GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx‎

Lines changed: 37 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -65,7 +65,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::prepare>(int32_t nBlock
6565
if (!tracks[i].OK()) {
6666
continue;
6767
}
68-
if (merger.Param().rec.tpc.dropSecondaryLegsInOutput && tracks[i].MergedLooper()) {
68+
if (tracks[i].MergedLooper()) {
6969
continue;
7070
}
7171

@@ -127,22 +127,23 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int32_t nBlocks
127127
for (int32_t iTmp = get_global_id(0); iTmp < nTracks; iTmp += get_global_size(0)) {
128128
TrackTPC oTrack;
129129
const int32_t i = trackSort[iTmp].x;
130-
auto snpIn = tracks[i].GetParam().GetSinPhi();
130+
const auto& track = tracks[i];
131+
auto snpIn = track.GetParam().GetSinPhi();
131132
if (snpIn > SNPThresh) {
132133
snpIn = SNPThresh;
133134
} else if (snpIn < -SNPThresh) {
134135
snpIn = -SNPThresh;
135136
}
136-
oTrack.set(tracks[i].GetParam().GetX(), tracks[i].GetAlpha(),
137-
{tracks[i].GetParam().GetY(), tracks[i].GetParam().GetZ(), snpIn, tracks[i].GetParam().GetDzDs(), tracks[i].GetParam().GetQPt()},
138-
{tracks[i].GetParam().GetCov(0),
139-
tracks[i].GetParam().GetCov(1), tracks[i].GetParam().GetCov(2),
140-
tracks[i].GetParam().GetCov(3), tracks[i].GetParam().GetCov(4), tracks[i].GetParam().GetCov(5),
141-
tracks[i].GetParam().GetCov(6), tracks[i].GetParam().GetCov(7), tracks[i].GetParam().GetCov(8), tracks[i].GetParam().GetCov(9),
142-
tracks[i].GetParam().GetCov(10), tracks[i].GetParam().GetCov(11), tracks[i].GetParam().GetCov(12), tracks[i].GetParam().GetCov(13), tracks[i].GetParam().GetCov(14)});
137+
oTrack.set(track.GetParam().GetX(), track.GetAlpha(),
138+
{track.GetParam().GetY(), track.GetParam().GetZ(), snpIn, track.GetParam().GetDzDs(), track.GetParam().GetQPt()},
139+
{track.GetParam().GetCov(0),
140+
track.GetParam().GetCov(1), track.GetParam().GetCov(2),
141+
track.GetParam().GetCov(3), track.GetParam().GetCov(4), track.GetParam().GetCov(5),
142+
track.GetParam().GetCov(6), track.GetParam().GetCov(7), track.GetParam().GetCov(8), track.GetParam().GetCov(9),
143+
track.GetParam().GetCov(10), track.GetParam().GetCov(11), track.GetParam().GetCov(12), track.GetParam().GetCov(13), track.GetParam().GetCov(14)});
143144

144-
oTrack.setChi2(tracks[i].GetParam().GetChi2());
145-
auto& outerPar = tracks[i].OuterParam();
145+
oTrack.setChi2(track.GetParam().GetChi2());
146+
auto& outerPar = track.OuterParam();
146147
if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) {
147148
if (param.dodEdxEnabled) {
148149
oTrack.setdEdx(tracksdEdx[i]);
@@ -189,13 +190,13 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int32_t nBlocks
189190
float t1 = 0, t2 = 0;
190191
int32_t sector1 = 0, sector2 = 0;
191192
const o2::tpc::ClusterNativeAccess* GPUrestrict() clusters = merger.GetConstantMem()->ioPtrs.clustersNative;
192-
for (uint32_t j = 0; j < tracks[i].NClusters(); j++) {
193-
if ((trackClusters[tracks[i].FirstClusterRef() + j].state & flagsReject) || (merger.ClusterAttachment()[trackClusters[tracks[i].FirstClusterRef() + j].num] & flagsRequired) != flagsRequired) {
193+
for (uint32_t j = 0; j < track.NClusters(); j++) {
194+
if ((trackClusters[track.FirstClusterRef() + j].state & flagsReject) || (merger.ClusterAttachment()[trackClusters[track.FirstClusterRef() + j].num] & flagsRequired) != flagsRequired) {
194195
continue;
195196
}
196-
int32_t clusterIdGlobal = trackClusters[tracks[i].FirstClusterRef() + j].num;
197-
int32_t sector = trackClusters[tracks[i].FirstClusterRef() + j].sector;
198-
int32_t globalRow = trackClusters[tracks[i].FirstClusterRef() + j].row;
197+
int32_t clusterIdGlobal = trackClusters[track.FirstClusterRef() + j].num;
198+
int32_t sector = trackClusters[track.FirstClusterRef() + j].sector;
199+
int32_t globalRow = trackClusters[track.FirstClusterRef() + j].row;
199200
int32_t clusterIdInRow = clusterIdGlobal - clusters->clusterOffset[sector][globalRow];
200201
clIndArr[nOutCl2] = clusterIdInRow;
201202
sectorIndexArr[nOutCl2] = sector;
@@ -210,25 +211,35 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int32_t nBlocks
210211
}
211212
}
212213

213-
bool cce = tracks[i].CCE() && ((sector1 < MAXSECTOR / 2) ^ (sector2 < MAXSECTOR / 2));
214+
if (track.PrevSegment() >= 0) {
215+
const GPUTPCGMMergedTrack* chkTrk = &tracks[track.PrevSegment()];
216+
while (chkTrk->PrevSegment() >= 0) {
217+
chkTrk = &tracks[chkTrk->PrevSegment()];
218+
}
219+
const auto& firstPrevCluster = trackClusters[chkTrk->FirstClusterRef()];
220+
t1 = clusters->clustersLinear[firstPrevCluster.num].getTime();
221+
sector1 = firstPrevCluster.sector;
222+
}
223+
224+
bool cce = track.CCE() && ((sector1 < MAXSECTOR / 2) ^ (sector2 < MAXSECTOR / 2));
214225
float time0 = 0.f, tFwd = 0.f, tBwd = 0.f;
215226
if (merger.Param().par.continuousTracking) {
216-
time0 = tracks[i].GetParam().GetTZOffset();
227+
time0 = track.GetParam().GetTZOffset();
217228
if (cce) {
218-
bool lastSide = trackClusters[tracks[i].FirstClusterRef()].sector < MAXSECTOR / 2;
229+
bool lastSide = trackClusters[track.FirstClusterRef()].sector < MAXSECTOR / 2;
219230
float delta = 0.f;
220-
for (uint32_t iCl = 1; iCl < tracks[i].NClusters(); iCl++) {
221-
auto& cacl1 = trackClusters[tracks[i].FirstClusterRef() + iCl];
231+
for (uint32_t iCl = 1; iCl < track.NClusters(); iCl++) {
232+
auto& cacl1 = trackClusters[track.FirstClusterRef() + iCl];
222233
if (lastSide ^ (cacl1.sector < MAXSECTOR / 2)) {
223234
auto& cl1 = clusters->clustersLinear[cacl1.num];
224-
auto& cl2 = clusters->clustersLinear[trackClusters[tracks[i].FirstClusterRef() + iCl - 1].num];
235+
auto& cl2 = clusters->clustersLinear[trackClusters[track.FirstClusterRef() + iCl - 1].num];
225236
delta = CAMath::Abs(cl1.getTime() - cl2.getTime()) * 0.5f;
226-
if (delta < MinDelta) {
227-
delta = MinDelta;
228-
}
229237
break;
230238
}
231239
}
240+
if (delta < MinDelta) {
241+
delta = MinDelta;
242+
}
232243
tFwd = tBwd = delta;
233244
} else {
234245
// estimate max/min time increments which still keep track in the physical limits of the TPC
@@ -261,7 +272,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int32_t nBlocks
261272
if (cce) {
262273
oTrack.setHasCSideClusters();
263274
oTrack.setHasASideClusters();
264-
} else if (tracks[i].CSide()) {
275+
} else if (track.CSide()) {
265276
oTrack.setHasCSideClusters();
266277
} else {
267278
oTrack.setHasASideClusters();

0 commit comments

Comments
 (0)