Skip to content

Commit 335d820

Browse files
davidrohrdavidrohr
committed
GPU TPC: Rename some variables with misleading name
1 parent 69911f8 commit 335d820

12 files changed

+48
-48
lines changed

‎GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAKernelsSpecialize.inc‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -109,13 +109,13 @@ inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed
109109
template <>
110110
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerSortTracks, 0>(const krnlSetupTime& _xyz)
111111
{
112-
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackOrderProcess(), processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
112+
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackOrderProcess(), processors()->tpcMerger.NMergedTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
113113
}
114114

115115
template <>
116116
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerSortTracksQPt, 0>(const krnlSetupTime& _xyz)
117117
{
118-
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackSort(), processors()->tpcMerger.NOutputTracks(), GPUTPCGMMergerSortTracksQPt_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
118+
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackSort(), processors()->tpcMerger.NMergedTracks(), GPUTPCGMMergerSortTracksQPt_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
119119
}
120120

121121
template <>

‎GPU/GPUTracking/DataCompression/GPUTPCCompression.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -125,7 +125,7 @@ void GPUTPCCompression::SetMaxData(const GPUTrackingInOutPointers& io)
125125
mMaxClusterFactorBase1024 = mMaxClusters > 100000000 ? mRec->MemoryScalers()->NTPCUnattachedHitsBase1024(mRec->GetParam().rec.tpc.rejectionStrategy) : 1024;
126126
mMaxClustersInCache = mMaxClusters * mMaxClusterFactorBase1024 / 1024;
127127
mMaxTrackClusters = mRec->GetConstantMem().tpcMerger.NOutputTrackClusters(); // TODO: Why is this not using ioPtrs? Could remove GPUConstantMem.h include
128-
mMaxTracks = mRec->GetConstantMem().tpcMerger.NOutputTracks();
128+
mMaxTracks = mRec->GetConstantMem().tpcMerger.NMergedTracks();
129129
if (mMaxClusters % 16) {
130130
mMaxClusters += 16 - (mMaxClusters % 16);
131131
}

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -201,7 +201,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
201201
const uint32_t iSector = iSectorRow / GPUCA_ROW_COUNT;
202202
const uint32_t iRow = iSectorRow % GPUCA_ROW_COUNT;
203203
const uint32_t idOffset = clusters->clusterOffset[iSector][iRow];
204-
const uint32_t idOffsetOut = clusters->clusterOffset[iSector][iRow] * compressor.mMaxClusterFactorBase1024 / 1024;
204+
const uint32_t idOffsetOut = clusters->clusterOffset[iSector][iRow] * compressor.mMaxClusterFactorBase1024 / 1024; // 32 bit enough for number of clusters per row * 1024
205205
const uint32_t idOffsetOutMax = ((const uint32_t*)clusters->clusterOffset[iSector])[iRow + 1] * compressor.mMaxClusterFactorBase1024 / 1024; // Array out of bounds access is ok, since it goes to the correct nClustersTotal
206206
if (iThread == nThreads - 1) {
207207
smem.nCount = 0;
@@ -214,7 +214,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
214214
const uint32_t nn = CAMath::nextMultipleOf<GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCompressionKernels_step1unattached)>(clusters->nClusters[iSector][iRow]);
215215
for (uint32_t i = iThread; i < nn + nThreads; i += nThreads) {
216216
const int32_t idx = idOffset + i;
217-
int32_t cidx = 0;
217+
int32_t storeCluster = 0;
218218
do {
219219
if (i >= clusters->nClusters[iSector][iRow]) {
220220
break;
@@ -239,13 +239,13 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
239239
break;
240240
}
241241
}
242-
cidx = 1;
242+
storeCluster = 1;
243243
} while (false);
244244

245245
GPUbarrier();
246-
int32_t myIndex = work_group_scan_inclusive_add(cidx);
246+
int32_t myIndex = work_group_scan_inclusive_add(storeCluster);
247247
int32_t storeLater = -1;
248-
if (cidx) {
248+
if (storeCluster) {
249249
if (smem.nCount + myIndex <= GPUCA_TPC_COMP_CHUNK_SIZE) {
250250
sortBuffer[smem.nCount + myIndex - 1] = i;
251251
} else {

‎GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -153,7 +153,7 @@ void GPUChainTracking::PrintMemoryStatistics()
153153
addToMap("TPC Sector TrackHits", usageMap, *processors()->tpcTrackers[i].NTrackHits(), processors()->tpcTrackers[i].NMaxTrackHits());
154154
}
155155
addToMap("TPC Clusterer Clusters", usageMap, mRec->MemoryScalers()->nTPCHits, mRec->MemoryScalers()->NTPCClusters(mRec->MemoryScalers()->nTPCdigits));
156-
addToMap("TPC Tracks", usageMap, processors()->tpcMerger.NOutputTracks(), processors()->tpcMerger.NMaxTracks());
156+
addToMap("TPC Tracks", usageMap, processors()->tpcMerger.NMergedTracks(), processors()->tpcMerger.NMaxTracks());
157157
addToMap("TPC TrackHits", usageMap, processors()->tpcMerger.NOutputTrackClusters(), processors()->tpcMerger.NMaxOutputTrackClusters());
158158

159159
if (mRec->GetProcessingSettings().createO2Output) {
@@ -181,7 +181,7 @@ void GPUChainTracking::PrintMemoryRelations()
181181
GPUInfo("MEMREL SectorTracks NCl %d NTrk %d", processors()->tpcTrackers[i].NHitsTotal(), *processors()->tpcTrackers[i].NTracks());
182182
GPUInfo("MEMREL SectorTrackHits NCl %d NTrkH %d", processors()->tpcTrackers[i].NHitsTotal(), *processors()->tpcTrackers[i].NTrackHits());
183183
}
184-
GPUInfo("MEMREL Tracks NCl %d NTrk %d", processors()->tpcMerger.NMaxClusters(), processors()->tpcMerger.NOutputTracks());
184+
GPUInfo("MEMREL Tracks NCl %d NTrk %d", processors()->tpcMerger.NMaxClusters(), processors()->tpcMerger.NMergedTracks());
185185
GPUInfo("MEMREL TrackHitss NCl %d NTrkH %d", processors()->tpcMerger.NMaxClusters(), processors()->tpcMerger.NOutputTrackClusters());
186186
}
187187

‎GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -220,7 +220,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
220220
mOutputQueue.clear();
221221
}
222222

223-
runKernel<GPUTPCGMMergerTrackFit>(doGPU ? GetGrid(Merger.NOutputTracks(), 0) : GetGridAuto(0), mergerSortTracks ? 1 : 0);
223+
runKernel<GPUTPCGMMergerTrackFit>(doGPU ? GetGrid(Merger.NMergedTracks(), 0) : GetGridAuto(0), mergerSortTracks ? 1 : 0);
224224
if (param().rec.tpc.retryRefit == 1) {
225225
runKernel<GPUTPCGMMergerTrackFit>(GetGridAuto(0), -1);
226226
}
@@ -233,7 +233,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
233233
runKernel<GPUTPCGMMergerFinalize, 1>(GetGridAuto(0, deviceType));
234234
runKernel<GPUTPCGMMergerFinalize, 2>(GetGridAuto(0, deviceType));
235235
if (param().rec.tpc.mergeLoopersAfterburner) {
236-
runKernel<GPUTPCGMMergerMergeLoopers, 0>(doGPU ? GetGrid(Merger.NOutputTracks(), 0, deviceType) : GetGridAuto(0, deviceType));
236+
runKernel<GPUTPCGMMergerMergeLoopers, 0>(doGPU ? GetGrid(Merger.NMergedTracks(), 0, deviceType) : GetGridAuto(0, deviceType));
237237
if (doGPU) {
238238
TransferMemoryResourceLinkToHost(RecoStep::TPCMerging, Merger.MemoryResMemory(), 0);
239239
SynchronizeStream(0); // TODO: could probably synchronize on an event after runKernel<GPUTPCGMMergerMergeLoopers, 1>
@@ -255,10 +255,10 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
255255
throw std::runtime_error("QA Scratch buffer exceeded");
256256
}
257257
}
258-
GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracks()), outputStream, 0, nullptr, waitEvent);
258+
GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NMergedTracks() * sizeof(*Merger.OutputTracks()), outputStream, 0, nullptr, waitEvent);
259259
waitEvent = nullptr;
260260
if (param().dodEdxEnabled) {
261-
GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NOutputTracks() * sizeof(*Merger.OutputTracksdEdx()), outputStream, 0);
261+
GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NMergedTracks() * sizeof(*Merger.OutputTracksdEdx()), outputStream, 0);
262262
}
263263
GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NOutputTrackClusters() * sizeof(*Merger.Clusters()), outputStream, 0);
264264
if (param().par.earlyTpcTransform) {
@@ -326,7 +326,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
326326
}
327327

328328
mIOPtrs.mergedTracks = Merger.OutputTracks();
329-
mIOPtrs.nMergedTracks = Merger.NOutputTracks();
329+
mIOPtrs.nMergedTracks = Merger.NMergedTracks();
330330
mIOPtrs.mergedTrackHits = Merger.Clusters();
331331
mIOPtrs.mergedTrackHitsXYZ = Merger.ClustersXYZ();
332332
mIOPtrs.nMergedTrackHits = Merger.NOutputTrackClusters();
@@ -340,7 +340,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
340340

341341
if (doGPU) {
342342
processorsShadow()->ioPtrs.mergedTracks = MergerShadow.OutputTracks();
343-
processorsShadow()->ioPtrs.nMergedTracks = Merger.NOutputTracks();
343+
processorsShadow()->ioPtrs.nMergedTracks = Merger.NMergedTracks();
344344
processorsShadow()->ioPtrs.mergedTrackHits = MergerShadow.Clusters();
345345
processorsShadow()->ioPtrs.mergedTrackHitsXYZ = MergerShadow.ClustersXYZ();
346346
processorsShadow()->ioPtrs.nMergedTrackHits = Merger.NOutputTrackClusters();

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -415,7 +415,7 @@ int32_t GPUTPCGMMerger::CheckSectors()
415415

416416
GPUd() void GPUTPCGMMerger::ClearTrackLinks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, bool output)
417417
{
418-
const int32_t n = output ? mMemory->nOutputTracks : SectorTrackInfoLocalTotal();
418+
const int32_t n = output ? mMemory->nMergedTracks : SectorTrackInfoLocalTotal();
419419
for (int32_t i = iBlock * nThreads + iThread; i < n; i += nThreads * nBlocks) {
420420
mTrackLinks[i] = -1;
421421
}
@@ -1271,7 +1271,7 @@ GPUd() void GPUTPCGMMerger::MergeCEFill(const GPUTPCGMSectorTrack* track, const
12711271
GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
12721272
{
12731273
const ClusterNative* cls = Param().par.earlyTpcTransform ? nullptr : mConstantMem->ioPtrs.clustersNative->clustersLinear;
1274-
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nThreads * nBlocks) {
1274+
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nThreads * nBlocks) {
12751275
if (mOutputTracks[i].CSide() == 0 && mTrackLinks[i] >= 0) {
12761276
if (mTrackLinks[mTrackLinks[i]] != (int32_t)i) {
12771277
continue;
@@ -1392,7 +1392,7 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i
13921392
}
13931393
}
13941394

1395-
// for (int32_t i = 0;i < mMemory->nOutputTracks;i++) {if (mOutputTracks[i].CCE() == false) {mOutputTracks[i].SetNClusters(0);mOutputTracks[i].SetOK(false);}} //Remove all non-CE tracks
1395+
// for (int32_t i = 0;i < mMemory->nMergedTracks;i++) {if (mOutputTracks[i].CCE() == false) {mOutputTracks[i].SetNClusters(0);mOutputTracks[i].SetOK(false);}} //Remove all non-CE tracks
13961396
}
13971397

13981398
namespace o2::gpu::internal
@@ -1533,7 +1533,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
15331533
nHits = 0;
15341534
for (int32_t ipart = 0; ipart < nParts; ipart++) {
15351535
const GPUTPCGMSectorTrack* t = trackParts[ipart];
1536-
CADEBUG(printf("Collect Track %d Part %d QPt %f DzDs %f\n", mMemory->nOutputTracks, ipart, t->QPt(), t->DzDs()));
1536+
CADEBUG(printf("Collect Track %d Part %d QPt %f DzDs %f\n", mMemory->nMergedTracks, ipart, t->QPt(), t->DzDs()));
15371537
int32_t nTrackHits = t->NClusters();
15381538
trackCluster* c2 = trackClusters + nHits + nTrackHits - 1;
15391539
for (int32_t i = 0; i < nTrackHits; i++, c2--) {
@@ -1678,10 +1678,10 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
16781678
cl[i].leg = trackClusters[i].leg;
16791679
}
16801680

1681-
uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nOutputTracks, 1u);
1681+
uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nMergedTracks, 1u);
16821682
if (iOutputTrack >= mNMaxTracks) {
16831683
raiseError(GPUErrors::ERROR_MERGER_TRACK_OVERFLOW, iOutputTrack, mNMaxTracks);
1684-
CAMath::AtomicExch(&mMemory->nOutputTracks, mNMaxTracks);
1684+
CAMath::AtomicExch(&mMemory->nMergedTracks, mNMaxTracks);
16851685
continue;
16861686
}
16871687

@@ -1718,9 +1718,9 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
17181718
p1.QPt() = 100.f / Param().rec.bz0Pt10MeV;
17191719
}
17201720

1721-
// if (nParts > 1) printf("Merged %d: QPt %f %d parts %d hits\n", mMemory->nOutputTracks, p1.QPt(), nParts, nHits);
1721+
// if (nParts > 1) printf("Merged %d: QPt %f %d parts %d hits\n", mMemory->nMergedTracks, p1.QPt(), nParts, nHits);
17221722

1723-
/*if (GPUQA::QAAvailable() && mRec->GetQA() && mRec->GetQA()->SuppressTrack(mMemory->nOutputTracks))
1723+
/*if (GPUQA::QAAvailable() && mRec->GetQA() && mRec->GetQA()->SuppressTrack(mMemory->nMergedTracks))
17241724
{
17251725
mergedTrack.SetOK(0);
17261726
mergedTrack.SetNClusters(0);
@@ -1742,14 +1742,14 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
17421742

17431743
GPUd() void GPUTPCGMMerger::SortTracksPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17441744
{
1745-
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nThreads * nBlocks) {
1745+
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nThreads * nBlocks) {
17461746
mTrackOrderProcess[i] = i;
17471747
}
17481748
}
17491749

17501750
GPUd() void GPUTPCGMMerger::PrepareClustersForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
17511751
{
1752-
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nBlocks * nThreads) {
1752+
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
17531753
mTrackSort[i] = i;
17541754
}
17551755
}
@@ -1784,7 +1784,7 @@ GPUd() void GPUTPCGMMerger::SortTracks(int32_t nBlocks, int32_t nThreads, int32_
17841784
) // clang-format on
17851785
};
17861786

1787-
GPUCommonAlgorithm::sortDeviceDynamic(mTrackOrderProcess, mTrackOrderProcess + mMemory->nOutputTracks, comp);
1787+
GPUCommonAlgorithm::sortDeviceDynamic(mTrackOrderProcess, mTrackOrderProcess + mMemory->nMergedTracks, comp);
17881788
#endif
17891789
}
17901790

@@ -1810,13 +1810,13 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int32_t nBlocks, int32_t nThreads, int
18101810
) // clang-format on
18111811
};
18121812

1813-
GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nOutputTracks, comp);
1813+
GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nMergedTracks, comp);
18141814
#endif
18151815
}
18161816

18171817
GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
18181818
{
1819-
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nBlocks * nThreads) {
1819+
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
18201820
mTrackOrderAttach[mTrackSort[i]] = i;
18211821
const GPUTPCGMMergedTrack& trk = mOutputTracks[i];
18221822
if (trk.OK()) {
@@ -1848,7 +1848,7 @@ GPUd() void GPUTPCGMMerger::PrepareClustersForFit2(int32_t nBlocks, int32_t nThr
18481848

18491849
GPUd() void GPUTPCGMMerger::Finalize0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
18501850
{
1851-
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nThreads * nBlocks) {
1851+
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nThreads * nBlocks) {
18521852
mTrackSort[mTrackOrderAttach[i]] = i;
18531853
}
18541854
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTrackClusters; i += nThreads * nBlocks) {
@@ -1858,7 +1858,7 @@ GPUd() void GPUTPCGMMerger::Finalize0(int32_t nBlocks, int32_t nThreads, int32_t
18581858

18591859
GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
18601860
{
1861-
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nOutputTracks; i += nThreads * nBlocks) {
1861+
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nThreads * nBlocks) {
18621862
const GPUTPCGMMergedTrack& trk = mOutputTracks[i];
18631863
if (!trk.OK() || trk.NClusters() == 0) {
18641864
continue;
@@ -1893,7 +1893,7 @@ GPUd() void GPUTPCGMMerger::Finalize2(int32_t nBlocks, int32_t nThreads, int32_t
18931893
GPUd() void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
18941894
{
18951895
const float lowPtThresh = Param().rec.tpc.rejectQPtB5 * 1.1f; // Might need to merge tracks above the threshold with parts below the threshold
1896-
for (uint32_t i = get_global_id(0); i < mMemory->nOutputTracks; i += get_global_size(0)) {
1896+
for (uint32_t i = get_global_id(0); i < mMemory->nMergedTracks; i += get_global_size(0)) {
18971897
const auto& trk = mOutputTracks[i];
18981898
const auto& p = trk.GetParam();
18991899
const float qptabs = CAMath::Abs(p.GetQPt());

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ class GPUTPCGMMerger : public GPUProcessor
6969
GPUAtomic(uint32_t) nRetryRefit;
7070
GPUAtomic(uint32_t) nLoopData;
7171
GPUAtomic(uint32_t) nUnpackedTracks;
72-
GPUAtomic(uint32_t) nOutputTracks;
72+
GPUAtomic(uint32_t) nMergedTracks;
7373
GPUAtomic(uint32_t) nOutputTrackClusters;
7474
GPUAtomic(uint32_t) nO2Tracks;
7575
GPUAtomic(uint32_t) nO2ClusRefs;
@@ -103,7 +103,7 @@ class GPUTPCGMMerger : public GPUProcessor
103103
void* SetPointersOutputState(void* mem);
104104
void* SetPointersMemory(void* mem);
105105

106-
GPUhdi() int32_t NOutputTracks() const { return mMemory->nOutputTracks; }
106+
GPUhdi() int32_t NMergedTracks() const { return mMemory->nMergedTracks; }
107107
GPUhdi() const GPUTPCGMMergedTrack* OutputTracks() const { return mOutputTracks; }
108108
GPUhdi() GPUTPCGMMergedTrack* OutputTracks() { return mOutputTracks; }
109109
GPUhdi() const GPUdEdxInfo* OutputTracksdEdx() const { return mOutputTracksdEdx; }

‎GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,7 @@ void GPUTPCGMMerger::DumpMergeRanges(std::ostream& out, int32_t withinSector, in
9494
void GPUTPCGMMerger::DumpTrackLinks(std::ostream& out, bool output, const char* type) const
9595
{
9696
out << "\nTPC Merger Links " << type << "\n";
97-
const int32_t n = output ? mMemory->nOutputTracks : SectorTrackInfoLocalTotal();
97+
const int32_t n = output ? mMemory->nMergedTracks : SectorTrackInfoLocalTotal();
9898
for (int32_t i = 0; i < n; i++) {
9999
if (mTrackLinks[i] != -1) {
100100
out << " " << i << ": " << mTrackLinks[i] << "\n";
@@ -138,7 +138,7 @@ void GPUTPCGMMerger::DumpCollected(std::ostream& out) const
138138
std::streamsize ss = out.precision();
139139
out << std::setprecision(2);
140140
out << "\nTPC Merger Collected Tracks\n";
141-
for (uint32_t i = 0; i < mMemory->nOutputTracks; i++) {
141+
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
142142
const auto& trk = mOutputTracks[i];
143143
const auto& p = trk.GetParam();
144144
out << " Track " << i << ": Loop " << trk.Looper() << " Alpha " << trk.GetAlpha() << " X " << p.GetX() << " offset " << p.GetTZOffset() << " Y " << p.GetY() << " Z " << p.GetZ() << " SPhi " << p.GetSinPhi() << " Tgl " << p.GetDzDs() << " QPt " << p.GetQPt() << " NCl " << trk.NClusters() << "\n";
@@ -150,7 +150,7 @@ void GPUTPCGMMerger::DumpMergeCE(std::ostream& out) const
150150
{
151151
DumpTrackLinks(out, true, " for CE merging");
152152
out << "\nTPC Merger Merge CE\n";
153-
for (uint32_t i = 0; i < mMemory->nOutputTracks; i++) {
153+
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
154154
const auto& trk = mOutputTracks[i];
155155
if (trk.CCE()) {
156156
out << " Track " << i << ": CCE\n";
@@ -162,11 +162,11 @@ void GPUTPCGMMerger::DumpFitPrepare(std::ostream& out) const
162162
{
163163
out << "\nTPC Merger Refit Prepare\n";
164164
out << " Sort\n";
165-
for (uint32_t i = 0; i < mMemory->nOutputTracks; i++) {
165+
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
166166
out << " " << i << ": " << mTrackOrderAttach[i] << "\n";
167167
}
168168
out << " Clusters\n";
169-
for (uint32_t j = 0; j < mMemory->nOutputTracks; j++) {
169+
for (uint32_t j = 0; j < mMemory->nMergedTracks; j++) {
170170
const auto& trk = mOutputTracks[j];
171171
out << " Track " << j << ": ";
172172
for (uint32_t i = trk.FirstClusterRef(); i < trk.FirstClusterRef() + trk.NClusters(); i++) {
@@ -195,7 +195,7 @@ void GPUTPCGMMerger::DumpRefit(std::ostream& out) const
195195
std::streamsize ss = out.precision();
196196
out << std::setprecision(2);
197197
out << "\nTPC Merger Refit\n";
198-
for (uint32_t i = 0; i < mMemory->nOutputTracks; i++) {
198+
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
199199
const auto& trk = mOutputTracks[i];
200200
if (trk.NClusters() == 0) {
201201
continue;
@@ -212,7 +212,7 @@ void GPUTPCGMMerger::DumpRefit(std::ostream& out) const
212212
void GPUTPCGMMerger::DumpFinal(std::ostream& out) const
213213
{
214214
out << "\nTPC Merger Finalized\n";
215-
for (uint32_t j = 0; j < mMemory->nOutputTracks; j++) {
215+
for (uint32_t j = 0; j < mMemory->nMergedTracks; j++) {
216216
const auto& trk = mOutputTracks[j];
217217
if (trk.NClusters() == 0) {
218218
continue;

0 commit comments

Comments
 (0)