Skip to content

Commit 1374b87

Browse files
davidrohrdavidrohr
committed
GPU TPC: Rename some kernels
1 parent c41507a commit 1374b87

File tree

10 files changed

+72
-71
lines changed

10 files changed

+72
-71
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: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -165,8 +165,8 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
165165
runKernel<GPUTPCGMMergerLinkExtrapolatedTracks>(GetGridAuto(0, deviceType));
166166
runKernel<GPUTPCGMMergerCollect>(GetGridAuto(0, deviceType));
167167
if (GetProcessingSettings().deterministicGPUReconstruction) {
168-
runKernel<GPUTPCGlobalDebugSortKernels, GPUTPCGlobalDebugSortKernels::extrapolatedTracks1>({{1, -WarpSize(), 0, deviceType}}, 1);
169-
runKernel<GPUTPCGlobalDebugSortKernels, GPUTPCGlobalDebugSortKernels::extrapolatedTracks2>({{1, -WarpSize(), 0, deviceType}}, 1);
168+
runKernel<GPUTPCGlobalDebugSortKernels, GPUTPCGlobalDebugSortKernels::mergedTracks1>({{1, -WarpSize(), 0, deviceType}}, 1);
169+
runKernel<GPUTPCGlobalDebugSortKernels, GPUTPCGlobalDebugSortKernels::mergedTracks2>({{1, -WarpSize(), 0, deviceType}}, 1);
170170
}
171171
DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingCollectedTracks, doGPU, Merger, &GPUTPCGMMerger::DumpCollected, *mDebugFile);
172172

@@ -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: 4 additions & 4 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;
@@ -1738,7 +1738,7 @@ GPUd() void GPUTPCGMMerger::SortTracksPrepare(int32_t nBlocks, int32_t nThreads,
17381738
}
17391739
}
17401740

1741-
GPUd() void GPUTPCGMMerger::PrepareClustersForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1741+
GPUd() void GPUTPCGMMerger::PrepareForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17421742
{
17431743
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
17441744
mTrackSort[i] = i;
@@ -1763,7 +1763,7 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int32_t nBlocks, int32_t nThreads, int
17631763
#endif
17641764
}
17651765

1766-
GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1766+
GPUd() void GPUTPCGMMerger::PrepareForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17671767
{
17681768
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
17691769
mTrackOrderAttach[mTrackSort[i]] = i;
@@ -1777,7 +1777,7 @@ GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int32_t nBlocks, int32_t nThr
17771777
}
17781778
}
17791779

1780-
GPUd() void GPUTPCGMMerger::PrepareClustersForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
1780+
GPUd() void GPUTPCGMMerger::PrepareForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17811781
{
17821782
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTrackClusters; i += nBlocks * nThreads) {
17831783
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: 26 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -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,25 @@ 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+
bool cce = track.CCE() && ((sector1 < MAXSECTOR / 2) ^ (sector2 < MAXSECTOR / 2));
214215
float time0 = 0.f, tFwd = 0.f, tBwd = 0.f;
215216
if (merger.Param().par.continuousTracking) {
216-
time0 = tracks[i].GetParam().GetTZOffset();
217+
time0 = track.GetParam().GetTZOffset();
217218
if (cce) {
218-
bool lastSide = trackClusters[tracks[i].FirstClusterRef()].sector < MAXSECTOR / 2;
219+
bool lastSide = trackClusters[track.FirstClusterRef()].sector < MAXSECTOR / 2;
219220
float delta = 0.f;
220-
for (uint32_t iCl = 1; iCl < tracks[i].NClusters(); iCl++) {
221-
auto& cacl1 = trackClusters[tracks[i].FirstClusterRef() + iCl];
221+
for (uint32_t iCl = 1; iCl < track.NClusters(); iCl++) {
222+
auto& cacl1 = trackClusters[track.FirstClusterRef() + iCl];
222223
if (lastSide ^ (cacl1.sector < MAXSECTOR / 2)) {
223224
auto& cl1 = clusters->clustersLinear[cacl1.num];
224-
auto& cl2 = clusters->clustersLinear[trackClusters[tracks[i].FirstClusterRef() + iCl - 1].num];
225+
auto& cl2 = clusters->clustersLinear[trackClusters[track.FirstClusterRef() + iCl - 1].num];
225226
delta = CAMath::Abs(cl1.getTime() - cl2.getTime()) * 0.5f;
226-
if (delta < MinDelta) {
227-
delta = MinDelta;
228-
}
229227
break;
230228
}
231229
}
230+
if (delta < MinDelta) {
231+
delta = MinDelta;
232+
}
232233
tFwd = tBwd = delta;
233234
} else {
234235
// estimate max/min time increments which still keep track in the physical limits of the TPC
@@ -261,7 +262,7 @@ GPUdii() void GPUTPCGMO2Output::Thread<GPUTPCGMO2Output::output>(int32_t nBlocks
261262
if (cce) {
262263
oTrack.setHasCSideClusters();
263264
oTrack.setHasASideClusters();
264-
} else if (tracks[i].CSide()) {
265+
} else if (track.CSide()) {
265266
oTrack.setHasCSideClusters();
266267
} else {
267268
oTrack.setHasASideClusters();

‎GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,7 @@ GPUdii() void GPUTPCGlobalDebugSortKernels::Thread<GPUTPCGlobalDebugSortKernels:
9494
}
9595

9696
template <>
97-
GPUdii() void GPUTPCGlobalDebugSortKernels::Thread<GPUTPCGlobalDebugSortKernels::extrapolatedTracks1>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter)
97+
GPUdii() void GPUTPCGlobalDebugSortKernels::Thread<GPUTPCGlobalDebugSortKernels::mergedTracks1>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter)
9898
{
9999
if (iThread || iBlock) {
100100
return;
@@ -112,7 +112,7 @@ GPUdii() void GPUTPCGlobalDebugSortKernels::Thread<GPUTPCGlobalDebugSortKernels:
112112
}
113113

114114
template <>
115-
GPUdii() void GPUTPCGlobalDebugSortKernels::Thread<GPUTPCGlobalDebugSortKernels::extrapolatedTracks2>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter)
115+
GPUdii() void GPUTPCGlobalDebugSortKernels::Thread<GPUTPCGlobalDebugSortKernels::mergedTracks2>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter)
116116
{
117117
if (iBlock) {
118118
return;

‎GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,8 @@ class GPUTPCGlobalDebugSortKernels : public GPUKernelTemplate
2929
enum K { defaultKernel = 0,
3030
clearIds = 0,
3131
sectorTracks = 1,
32-
extrapolatedTracks1 = 2,
33-
extrapolatedTracks2 = 3,
32+
mergedTracks1 = 2,
33+
mergedTracks2 = 3,
3434
borderTracks = 4 };
3535
GPUhdi() constexpr static GPUDataTypes::RecoStep GetRecoStep() { return GPUDataTypes::RecoStep::TPCMerging; }
3636
typedef GPUTPCGMMerger processorType;

0 commit comments

Comments
 (0)