Skip to content

Commit b8a72f7

Browse files
davidrohrdavidrohr
committed
GPU TPC: Order legs in descending way and store leg id per track not cluster
1 parent 636e88c commit b8a72f7

File tree

12 files changed

+52
-87
lines changed

12 files changed

+52
-87
lines changed

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 1 addition & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,6 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
3232
GPUTPCCompression& GPUrestrict() compressor = processors.tpcCompressor;
3333
const GPUParam& GPUrestrict() param = processors.param;
3434

35-
uint8_t lastLeg = 0;
3635
int32_t myTrack = 0;
3736
for (uint32_t i = get_global_id(0); i < ioPtrs.nMergedTracks; i += get_global_size(0)) {
3837
GPUbarrierWarp();
@@ -75,9 +74,6 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
7574
if ((hit.sector < GPUCA_NSECTORS) ^ (lastSector < GPUCA_NSECTORS)) {
7675
break;
7776
}
78-
if (lastLeg != hit.leg && track.Mirror()) {
79-
break;
80-
}
8177
if (track.Propagate(geo.Row2X(hit.row), param.SectorParam[hit.sector].Alpha)) {
8278
break;
8379
}
@@ -93,7 +89,6 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
9389

9490
myTrack = CAMath::AtomicAdd(&compressor.mMemory->nStoredTracks, 1u);
9591
compressor.mAttachedClusterFirstIndex[myTrack] = trk.FirstClusterRef();
96-
lastLeg = hit.leg;
9792
c.qPtA[myTrack] = qpt;
9893
c.rowA[myTrack] = hit.row;
9994
c.sliceA[myTrack] = hit.sector;
@@ -114,12 +109,11 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
114109
sector -= lastSector;
115110
}
116111
c.rowDiffA[cidx] = row;
117-
c.sliceLegDiffA[cidx] = (hit.leg == lastLeg ? 0 : compressor.NSECTORS) + sector;
112+
c.sliceLegDiffA[cidx] = sector;
118113
float pad = CAMath::Max(0.f, CAMath::Min((float)geo.NPads(GPUCA_ROW_COUNT - 1), track.LinearY2Pad(hit.sector, track.Y(), geo.PadWidth(hit.row), geo.NPads(hit.row))));
119114
c.padResA[cidx] = orgCl.padPacked - orgCl.packPad(pad);
120115
float time = CAMath::Max(0.f, geo.LinearZ2Time(hit.sector, track.Z() + zOffset));
121116
c.timeResA[cidx] = (orgCl.getTimePacked() - orgCl.packTime(time)) & 0xFFFFFF;
122-
lastLeg = hit.leg;
123117
}
124118
uint16_t qtot = orgCl.qTot, qmax = orgCl.qMax;
125119
uint8_t sigmapad = orgCl.sigmaPadPacked, sigmatime = orgCl.sigmaTimePacked;

‎GPU/GPUTracking/DataTypes/GPUTPCGMMergedTrackHit.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,7 @@ namespace o2::gpu
2121
{
2222
struct GPUTPCGMMergedTrackHit {
2323
uint32_t num;
24-
uint8_t sector, row, leg, state;
24+
uint8_t sector, row, state;
2525

2626
// NOTE: the lower states must match those from ClusterNative!
2727
// TODO: take them directly from clusterNative header.

‎GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -49,6 +49,7 @@ class GPUTPCGMMergedTrack
4949
GPUd() bool MergedLooperConnected() const { return mFlags & 0x20; }
5050
GPUd() bool MergedLooper() const { return mFlags & 0x30; }
5151
GPUd() int32_t PrevSegment() const { return mPrevSegment; }
52+
GPUd() uint8_t Leg() const { return mLeg; }
5253
GPUd() uint8_t Flags() const { return mFlags; }
5354

5455
GPUd() void SetNClusters(int32_t v) { mNClusters = v; }
@@ -57,6 +58,7 @@ class GPUTPCGMMergedTrack
5758
GPUd() void SetParam(const GPUTPCGMTrackParam& v) { mParam = v; }
5859
GPUd() void SetAlpha(float v) { mAlpha = v; }
5960
GPUd() void SetPrevSegment(int32_t v) { mPrevSegment = v; }
61+
GPUd() void SetLeg(uint8_t v) { mLeg = v; }
6062
GPUd() void SetOK(bool v)
6163
{
6264
if (v) {
@@ -121,6 +123,7 @@ class GPUTPCGMMergedTrack
121123
uint32_t mNClusters; //* number of track clusters
122124
uint32_t mNClustersFitted; //* number of clusters used in fit
123125
uint8_t mFlags;
126+
uint8_t mLeg;
124127

125128
#if !defined(GPUCA_STANDALONE)
126129
ClassDefNV(GPUTPCGMMergedTrack, 0);

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 11 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -179,7 +179,6 @@ void GPUTPCGMMerger::CheckMergedTracks()
179179
if (track.PrevNeighbour() >= 0) {
180180
continue;
181181
}
182-
int32_t leg = 0;
183182
GPUTPCGMSectorTrack *trbase = &track, *tr = &track;
184183
while (true) {
185184
int32_t iTrk = tr - mSectorTrackInfos;
@@ -200,7 +199,6 @@ void GPUTPCGMMerger::CheckMergedTracks()
200199
if (tr->PrevSegmentNeighbour() >= 0) {
201200
break;
202201
}
203-
leg++;
204202
continue;
205203
}
206204
break;
@@ -1463,15 +1461,15 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
14631461
int32_t leg = 0;
14641462
int32_t lastMergedSegment = -1;
14651463
while (true) {
1466-
if (trbase) {
1464+
if (trbase && !Param().rec.tpc.dropLoopers) {
14671465
int32_t jtr = trbase->NextNeighbour();
14681466
if (jtr >= 0) {
14691467
trbase = &(mSectorTrackInfos[jtr]);
14701468
if (trbase->PrevSegmentNeighbour() >= 0) {
14711469
trbase = nullptr;
14721470
} else {
14731471
trbase->SetPrevSegmentNeighbour(1000000001);
1474-
leg++;
1472+
leg--;
14751473
}
14761474
} else {
14771475
trbase = nullptr;
@@ -1492,7 +1490,12 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
14921490
}
14931491
itr += nThreads * nBlocks;
14941492
trbase->SetPrevSegmentNeighbour(1000000000);
1493+
int32_t jtr = trbase->NextNeighbour();
14951494
leg = 0;
1495+
while (jtr >= 0) {
1496+
leg++;
1497+
jtr = mSectorTrackInfos[jtr].NextNeighbour();
1498+
}
14961499
lastMergedSegment = -1;
14971500
}
14981501

@@ -1535,7 +1538,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
15351538
}
15361539

15371540
// unpack and sort clusters
1538-
if (nParts > 1 && leg == 0) {
1541+
if (nParts > 1) {
15391542
GPUCommonAlgorithm::sort(trackParts, trackParts + nParts, [](const GPUTPCGMSectorTrack* a, const GPUTPCGMSectorTrack* b) {
15401543
GPUCA_DETERMINISTIC_CODE( // clang-format off
15411544
if (a->X() != b->X()) {
@@ -1554,11 +1557,6 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
15541557
});
15551558
}
15561559

1557-
if (Param().rec.tpc.dropLoopers && leg > 0) {
1558-
nParts = 1;
1559-
leg = 0;
1560-
}
1561-
15621560
trackCluster trackClusters[kMaxClusters];
15631561
nHits = 0;
15641562
for (int32_t ipart = 0; ipart < nParts; ipart++) {
@@ -1570,7 +1568,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
15701568
const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()];
15711569
const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i];
15721570
uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0];
1573-
*c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector(), t->Leg()};
1571+
*c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector()};
15741572
}
15751573
nHits += nTrackHits;
15761574
}
@@ -1651,7 +1649,6 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
16511649
cl[i].row = trackClusters[i].row;
16521650
cl[i].num = trackClusters[i].id;
16531651
cl[i].sector = trackClusters[i].sector;
1654-
cl[i].leg = trackClusters[i].leg;
16551652
}
16561653

16571654
uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nMergedTracks, 1u);
@@ -1664,7 +1661,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
16641661
GPUTPCGMMergedTrack& mergedTrack = mMergedTracks[iOutputTrack];
16651662
mergedTrack.SetFlags(0);
16661663
mergedTrack.SetOK(true);
1667-
mergedTrack.SetLooper(leg > 0 || trbase->NextNeighbour() >= 0);
1664+
mergedTrack.SetLooper(leg > 0 || lastMergedSegment >= 0);
16681665
mergedTrack.SetNClusters(nHits);
16691666
mergedTrack.SetFirstClusterRef(iMergedTrackFirstCluster);
16701667
GPUTPCGMTrackParam& p1 = mergedTrack.Param();
@@ -1799,7 +1796,6 @@ GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t
17991796
if (!trk.OK() || trk.NClusters() == 0) {
18001797
continue;
18011798
}
1802-
uint8_t goodLeg = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].leg;
18031799
for (uint32_t j = 0; j < trk.NClusters(); j++) {
18041800
int32_t id = mClusters[trk.FirstClusterRef() + j].num;
18051801
uint32_t weight = mTrackOrderAttach[i] | attachAttached;
@@ -1809,7 +1805,7 @@ GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t
18091805
} else if (clusterState & GPUTPCGMMergedTrackHit::flagHighIncl) {
18101806
weight |= attachHighIncl;
18111807
}
1812-
if (mClusters[trk.FirstClusterRef() + j].leg == goodLeg) {
1808+
if (trk.Leg() == 0) {
18131809
weight |= attachGoodLeg;
18141810
}
18151811
CAMath::AtomicMax(&mClusterAttachment[id], weight);

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -82,7 +82,6 @@ class GPUTPCGMMerger : public GPUProcessor
8282
uint32_t id;
8383
uint8_t row;
8484
uint8_t sector;
85-
uint8_t leg;
8685
};
8786

8887
struct tmpSort {

‎GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx‎

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -65,14 +65,15 @@ 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()) {
69+
continue;
70+
}
71+
6872
uint32_t nCl = 0;
6973
for (uint32_t j = 0; j < tracks[i].NClusters(); j++) {
7074
if ((trackClusters[tracks[i].FirstClusterRef() + j].state & flagsReject) || (merger.ClusterAttachment()[trackClusters[tracks[i].FirstClusterRef() + j].num] & flagsRequired) != flagsRequired) {
7175
continue;
7276
}
73-
if (merger.Param().rec.tpc.dropSecondaryLegsInOutput && trackClusters[tracks[i].FirstClusterRef() + j].leg != trackClusters[tracks[i].FirstClusterRef() + tracks[i].NClusters() - 1].leg) {
74-
continue;
75-
}
7677
nCl++;
7778
}
7879
if (nCl == 0) {
@@ -192,9 +193,6 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int32_t nBlocks
192193
if ((trackClusters[tracks[i].FirstClusterRef() + j].state & flagsReject) || (merger.ClusterAttachment()[trackClusters[tracks[i].FirstClusterRef() + j].num] & flagsRequired) != flagsRequired) {
193194
continue;
194195
}
195-
if (merger.Param().rec.tpc.dropSecondaryLegsInOutput && trackClusters[tracks[i].FirstClusterRef() + j].leg != trackClusters[tracks[i].FirstClusterRef() + tracks[i].NClusters() - 1].leg) {
196-
continue;
197-
}
198196
int32_t clusterIdGlobal = trackClusters[tracks[i].FirstClusterRef() + j].num;
199197
int32_t sector = trackClusters[tracks[i].FirstClusterRef() + j].sector;
200198
int32_t globalRow = trackClusters[tracks[i].FirstClusterRef() + j].row;

0 commit comments

Comments
 (0)