Skip to content

Commit 747fb86

Browse files
davidrohrdavidrohr
committed
GPU TPC: Some more member variable renaming
1 parent f75693d commit 747fb86

10 files changed

+59
-59
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,13 +27,13 @@ inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed
2727
template <>
2828
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerSortTracks, 0>(const krnlSetupTime& _xyz)
2929
{
30-
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackOrderProcess(), processors()->tpcMerger.NMergedTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
30+
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackOrderProcess(), processors()->tpcMerger.NMergedTracks(), GPUTPCGMMergerSortTracks_comp(mProcessorsShadow->tpcMerger.MergedTracks()));
3131
}
3232

3333
template <>
3434
inline void GPUCA_M_CAT(GPUReconstruction, GPUCA_GPUTYPE)::runKernelBackendTimed<GPUTPCGMMergerSortTracksQPt, 0>(const krnlSetupTime& _xyz)
3535
{
36-
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackSort(), processors()->tpcMerger.NMergedTracks(), GPUTPCGMMergerSortTracksQPt_comp(mProcessorsShadow->tpcMerger.OutputTracks()));
36+
GPUCommonAlgorithm::sortOnDevice(this, _xyz.x.stream, mProcessorsShadow->tpcMerger.TrackSort(), processors()->tpcMerger.NMergedTracks(), GPUTPCGMMergerSortTracksQPt_comp(mProcessorsShadow->tpcMerger.MergedTracks()));
3737
}
3838

3939
template <>

‎GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -256,10 +256,10 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
256256
throw std::runtime_error("QA Scratch buffer exceeded");
257257
}
258258
}
259-
GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracks(), MergerShadowAll.OutputTracks(), Merger.NMergedTracks() * sizeof(*Merger.OutputTracks()), outputStream, 0, nullptr, waitEvent);
259+
GPUMemCpy(RecoStep::TPCMerging, Merger.MergedTracks(), MergerShadowAll.MergedTracks(), Merger.NMergedTracks() * sizeof(*Merger.MergedTracks()), outputStream, 0, nullptr, waitEvent);
260260
waitEvent = nullptr;
261261
if (param().dodEdxEnabled) {
262-
GPUMemCpy(RecoStep::TPCMerging, Merger.OutputTracksdEdx(), MergerShadowAll.OutputTracksdEdx(), Merger.NMergedTracks() * sizeof(*Merger.OutputTracksdEdx()), outputStream, 0);
262+
GPUMemCpy(RecoStep::TPCMerging, Merger.MergedTracksdEdx(), MergerShadowAll.MergedTracksdEdx(), Merger.NMergedTracks() * sizeof(*Merger.MergedTracksdEdx()), outputStream, 0);
263263
}
264264
GPUMemCpy(RecoStep::TPCMerging, Merger.Clusters(), MergerShadowAll.Clusters(), Merger.NOutputTrackClusters() * sizeof(*Merger.Clusters()), outputStream, 0);
265265
if (param().par.earlyTpcTransform) {
@@ -326,7 +326,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput)
326326
mRec->ReturnVolatileDeviceMemory();
327327
}
328328

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

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

‎GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ int32_t GPUChainTracking::RunRefit()
3333
SetupGPUProcessor(&Refit, false);
3434
RefitShadow.SetPtrsFromGPUConstantMem(processorsShadow(), doGPU ? &processorsDevice()->param : nullptr);
3535
RefitShadow.SetPropagator(doGPU ? processorsShadow()->calibObjects.o2Propagator : GetO2Propagator());
36-
RefitShadow.mPTracks = (doGPU ? processorsShadow() : processors())->tpcMerger.OutputTracks();
36+
RefitShadow.mPTracks = (doGPU ? processorsShadow() : processors())->tpcMerger.MergedTracks();
3737
WriteToConstantMemory(RecoStep::Refit, (char*)&processors()->trackingRefit - (char*)processors(), &RefitShadow, sizeof(RefitShadow), 0);
3838
// TransferMemoryResourcesToGPU(RecoStep::Refit, &Refit, 0);
3939
if (param().rec.trackingRefitGPUModel) {

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx‎

Lines changed: 19 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -365,11 +365,11 @@ void* GPUTPCGMMerger::SetPointersRefitScratch(void* mem)
365365

366366
void* GPUTPCGMMerger::SetPointersOutput(void* mem)
367367
{
368-
computePointerWithAlignment(mem, mOutputTracks, mNMaxTracks);
368+
computePointerWithAlignment(mem, mMergedTracks, mNMaxTracks);
369369
if (mRec->GetParam().dodEdxEnabled) {
370-
computePointerWithAlignment(mem, mOutputTracksdEdx, mNMaxTracks);
370+
computePointerWithAlignment(mem, mMergedTracksdEdx, mNMaxTracks);
371371
if (mRec->GetParam().rec.tpc.dEdxClusterRejectionFlagMask != mRec->GetParam().rec.tpc.dEdxClusterRejectionFlagMaskAlt) {
372-
computePointerWithAlignment(mem, mOutputTracksdEdxAlt, mNMaxTracks);
372+
computePointerWithAlignment(mem, mMergedTracksdEdxAlt, mNMaxTracks);
373373
}
374374
}
375375
computePointerWithAlignment(mem, mClusters, mNMaxOutputTrackClusters);
@@ -1318,7 +1318,7 @@ GPUd() void GPUTPCGMMerger::MergeCEFill(const GPUTPCGMSectorTrack* track, const
13181318
const float x0 = GPUTPCGeometry::Row2X(attempt == 0 ? 63 : cls.row);
13191319
if (track->TransportToX(this, x0, Param().bzCLight, b, GPUCA_MAX_SIN_PHI_LOW)) {
13201320
b.SetTrackID(itr);
1321-
b.SetNClusters(mOutputTracks[itr].NClusters());
1321+
b.SetNClusters(mMergedTracks[itr].NClusters());
13221322
if (CAMath::Abs(b.Cov()[4]) >= 0.5f) {
13231323
b.SetCov(4, 0.5f); // TODO: Is this needed and better than the cut in BorderTrack?
13241324
}
@@ -1339,11 +1339,11 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i
13391339
{
13401340
const ClusterNative* cls = Param().par.earlyTpcTransform ? nullptr : mConstantMem->ioPtrs.clustersNative->clustersLinear;
13411341
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nThreads * nBlocks) {
1342-
if (mOutputTracks[i].CSide() == 0 && mTrackLinks[i] >= 0) {
1342+
if (mMergedTracks[i].CSide() == 0 && mTrackLinks[i] >= 0) {
13431343
if (mTrackLinks[mTrackLinks[i]] != (int32_t)i) {
13441344
continue;
13451345
}
1346-
GPUTPCGMMergedTrack* trk[2] = {&mOutputTracks[i], &mOutputTracks[mTrackLinks[i]]};
1346+
GPUTPCGMMergedTrack* trk[2] = {&mMergedTracks[i], &mMergedTracks[mTrackLinks[i]]};
13471347

13481348
if (!trk[1]->OK() || trk[1]->CCE()) {
13491349
continue;
@@ -1459,7 +1459,7 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i
14591459
}
14601460
}
14611461

1462-
// 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
1462+
// for (int32_t i = 0;i < mMemory->nMergedTracks;i++) {if (mMergedTracks[i].CCE() == false) {mMergedTracks[i].SetNClusters(0);mMergedTracks[i].SetOK(false);}} //Remove all non-CE tracks
14631463
}
14641464

14651465
namespace o2::gpu::internal
@@ -1752,7 +1752,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
17521752
continue;
17531753
}
17541754

1755-
GPUTPCGMMergedTrack& mergedTrack = mOutputTracks[iOutputTrack];
1755+
GPUTPCGMMergedTrack& mergedTrack = mMergedTracks[iOutputTrack];
17561756

17571757
mergedTrack.SetFlags(0);
17581758
mergedTrack.SetOK(1);
@@ -1825,7 +1825,7 @@ GPUd() void GPUTPCGMMerger::SortTracks(int32_t nBlocks, int32_t nThreads, int32_
18251825
{
18261826
#ifndef GPUCA_SPECIALIZE_THRUST_SORTS
18271827
if (iThread == 0 && iBlock == 0) {
1828-
GPUCommonAlgorithm::sortDeviceDynamic(mTrackOrderProcess, mTrackOrderProcess + mMemory->nMergedTracks, GPUTPCGMMergerSortTracks_comp(mOutputTracks));
1828+
GPUCommonAlgorithm::sortDeviceDynamic(mTrackOrderProcess, mTrackOrderProcess + mMemory->nMergedTracks, GPUTPCGMMergerSortTracks_comp(mMergedTracks));
18291829
}
18301830
#endif
18311831
}
@@ -1834,7 +1834,7 @@ GPUd() void GPUTPCGMMerger::SortTracksQPt(int32_t nBlocks, int32_t nThreads, int
18341834
{
18351835
#ifndef GPUCA_SPECIALIZE_THRUST_SORTS
18361836
if (iThread == 0 && iBlock == 0) {
1837-
GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nMergedTracks, GPUTPCGMMergerSortTracksQPt_comp(mOutputTracks));
1837+
GPUCommonAlgorithm::sortDeviceDynamic(mTrackSort, mTrackSort + mMemory->nMergedTracks, GPUTPCGMMergerSortTracksQPt_comp(mMergedTracks));
18381838
}
18391839
#endif
18401840
}
@@ -1843,7 +1843,7 @@ GPUd() void GPUTPCGMMerger::PrepareClustersForFit1(int32_t nBlocks, int32_t nThr
18431843
{
18441844
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nBlocks * nThreads) {
18451845
mTrackOrderAttach[mTrackSort[i]] = i;
1846-
const GPUTPCGMMergedTrack& trk = mOutputTracks[i];
1846+
const GPUTPCGMMergedTrack& trk = mMergedTracks[i];
18471847
if (trk.OK()) {
18481848
for (uint32_t j = 0; j < trk.NClusters(); j++) {
18491849
mClusterAttachment[mClusters[trk.FirstClusterRef() + j].num] = attachAttached | attachGood;
@@ -1884,7 +1884,7 @@ GPUd() void GPUTPCGMMerger::Finalize0(int32_t nBlocks, int32_t nThreads, int32_t
18841884
GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
18851885
{
18861886
for (uint32_t i = iBlock * nThreads + iThread; i < mMemory->nMergedTracks; i += nThreads * nBlocks) {
1887-
const GPUTPCGMMergedTrack& trk = mOutputTracks[i];
1887+
const GPUTPCGMMergedTrack& trk = mMergedTracks[i];
18881888
if (!trk.OK() || trk.NClusters() == 0) {
18891889
continue;
18901890
}
@@ -1919,7 +1919,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads,
19191919
{
19201920
const float lowPtThresh = Param().rec.tpc.rejectQPtB5 * 1.1f; // Might need to merge tracks above the threshold with parts below the threshold
19211921
for (uint32_t i = get_global_id(0); i < mMemory->nMergedTracks; i += get_global_size(0)) {
1922-
const auto& trk = mOutputTracks[i];
1922+
const auto& trk = mMergedTracks[i];
19231923
const auto& p = trk.GetParam();
19241924
const float qptabs = CAMath::Abs(p.GetQPt());
19251925
if (trk.NClusters() && qptabs * Param().qptB5Scaler > 5.f && qptabs * Param().qptB5Scaler <= lowPtThresh) {
@@ -1983,7 +1983,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads,
19831983
#if GPUCA_MERGE_LOOPER_MC && !defined(GPUCA_GPUCODE)
19841984
std::vector<int64_t> paramLabels(mMemory->nLooperMatchCandidates);
19851985
for (uint32_t i = 0; i < mMemory->nLooperMatchCandidates; i++) {
1986-
paramLabels[i] = GetTrackLabel(mOutputTracks[params[i].id]);
1986+
paramLabels[i] = GetTrackLabel(mMergedTracks[params[i].id]);
19871987
}
19881988
/*std::vector<bool> dropped(mMemory->nLooperMatchCandidates);
19891989
std::vector<bool> droppedMC(mMemory->nLooperMatchCandidates);
@@ -2005,8 +2005,8 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads,
20052005
// bs |= 1;
20062006
continue;
20072007
}
2008-
const auto& trk1 = mOutputTracks[params[i].id];
2009-
const auto& trk2 = mOutputTracks[params[j].id];
2008+
const auto& trk1 = mMergedTracks[params[i].id];
2009+
const auto& trk2 = mMergedTracks[params[j].id];
20102010
const auto& param1 = trk1.GetParam();
20112011
const auto& param2 = trk2.GetParam();
20122012
if (CAMath::Abs(param1.GetDzDs()) > 0.03f && CAMath::Abs(param2.GetDzDs()) > 0.03f && param1.GetDzDs() * param2.GetDzDs() * param1.GetQPt() * param2.GetQPt() < 0) {
@@ -2045,7 +2045,7 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads,
20452045
const int64_t label2 = paramLabels[j];
20462046
bool labelEQ = label1 != -1 && label1 == label2;
20472047
if (1 || EQ || labelEQ) {
2048-
// printf("Matching track %d/%d %u-%u (%ld/%ld): dist %f side %d %d, tgl %f %f, qpt %f %f, x %f %f, y %f %f\n", (int32_t)EQ, (int32_t)labelEQ, i, j, label1, label2, d, (int32_t)mOutputTracks[params[i].id].CSide(), (int32_t)mOutputTracks[params[j].id].CSide(), params[i].tgl, params[j].tgl, params[i].qpt, params[j].qpt, params[i].x, params[j].x, params[i].y, params[j].y);
2048+
// printf("Matching track %d/%d %u-%u (%ld/%ld): dist %f side %d %d, tgl %f %f, qpt %f %f, x %f %f, y %f %f\n", (int32_t)EQ, (int32_t)labelEQ, i, j, label1, label2, d, (int32_t)mMergedTracks[params[i].id].CSide(), (int32_t)mMergedTracks[params[j].id].CSide(), params[i].tgl, params[j].tgl, params[i].qpt, params[j].qpt, params[i].x, params[j].x, params[i].y, params[j].y);
20492049
static auto& tup = GPUROOTDump<TNtuple>::get("mergeloopers", "labeleq:sides:d2xy:tgl1:tgl2:qpt1:qpt2:dz:dzcorr:dtgl:dqpt:dznorm:bs");
20502050
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);
20512051
static auto tup2 = GPUROOTDump<TNtuple>::getNew("mergeloopers2", "labeleq:refz1:refz2:tgl1:tgl2:qpt1:qpt2:snp1:snp2:a1:a2:dzn:phasecor:phasedir:dzcorr");
@@ -2063,9 +2063,9 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads,
20632063
}*/
20642064
#endif
20652065
if (EQ) {
2066-
mOutputTracks[params[j].id].SetMergedLooper(true);
2066+
mMergedTracks[params[j].id].SetMergedLooper(true);
20672067
if (CAMath::Abs(param2.GetQPt() * Param().qptB5Scaler) >= Param().rec.tpc.rejectQPtB5) {
2068-
mOutputTracks[params[i].id].SetMergedLooper(true);
2068+
mMergedTracks[params[i].id].SetMergedLooper(true);
20692069
}
20702070
}
20712071
}

‎GPU/GPUTracking/Merger/GPUTPCGMMerger.h‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -104,12 +104,12 @@ class GPUTPCGMMerger : public GPUProcessor
104104
void* SetPointersMemory(void* mem);
105105

106106
GPUhdi() int32_t NMergedTracks() const { return mMemory->nMergedTracks; }
107-
GPUhdi() const GPUTPCGMMergedTrack* OutputTracks() const { return mOutputTracks; }
108-
GPUhdi() GPUTPCGMMergedTrack* OutputTracks() { return mOutputTracks; }
109-
GPUhdi() const GPUdEdxInfo* OutputTracksdEdx() const { return mOutputTracksdEdx; }
110-
GPUhdi() GPUdEdxInfo* OutputTracksdEdx() { return mOutputTracksdEdx; }
111-
GPUhdi() const GPUdEdxInfo* OutputTracksdEdxAlt() const { return mOutputTracksdEdxAlt; }
112-
GPUhdi() GPUdEdxInfo* OutputTracksdEdxAlt() { return mOutputTracksdEdxAlt; }
107+
GPUhdi() const GPUTPCGMMergedTrack* MergedTracks() const { return mMergedTracks; }
108+
GPUhdi() GPUTPCGMMergedTrack* MergedTracks() { return mMergedTracks; }
109+
GPUhdi() const GPUdEdxInfo* MergedTracksdEdx() const { return mMergedTracksdEdx; }
110+
GPUhdi() GPUdEdxInfo* MergedTracksdEdx() { return mMergedTracksdEdx; }
111+
GPUhdi() const GPUdEdxInfo* MergedTracksdEdxAlt() const { return mMergedTracksdEdxAlt; }
112+
GPUhdi() GPUdEdxInfo* MergedTracksdEdxAlt() { return mMergedTracksdEdxAlt; }
113113
GPUhdi() uint32_t NClusters() const { return mNClusters; }
114114
GPUhdi() uint32_t NMaxClusters() const { return mNMaxClusters; }
115115
GPUhdi() uint32_t NMaxTracks() const { return mNMaxTracks; }
@@ -262,9 +262,9 @@ class GPUTPCGMMerger : public GPUProcessor
262262
uint16_t mMemoryResOutputO2Scratch = (uint16_t)-1;
263263

264264
int32_t mNClusters = 0; // Total number of incoming clusters (from sector tracks)
265-
GPUTPCGMMergedTrack* mOutputTracks = nullptr; //* array of output merged tracks
266-
GPUdEdxInfo* mOutputTracksdEdx = nullptr; //* dEdx information
267-
GPUdEdxInfo* mOutputTracksdEdxAlt = nullptr; //* dEdx alternative information
265+
GPUTPCGMMergedTrack* mMergedTracks = nullptr; //* array of output merged tracks
266+
GPUdEdxInfo* mMergedTracksdEdx = nullptr; //* dEdx information
267+
GPUdEdxInfo* mMergedTracksdEdxAlt = nullptr; //* dEdx alternative information
268268
GPUTPCGMSectorTrack* mSectorTrackInfos = nullptr; //* additional information for sector tracks
269269
int32_t* mSectorTrackInfoIndex = nullptr;
270270
GPUTPCGMMergedTrackHit* mClusters = nullptr;

‎GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -139,7 +139,7 @@ void GPUTPCGMMerger::DumpCollected(std::ostream& out) const
139139
out << std::setprecision(2);
140140
out << "\nTPC Merger Collected Tracks\n";
141141
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
142-
const auto& trk = mOutputTracks[i];
142+
const auto& trk = mMergedTracks[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";
145145
}
@@ -151,7 +151,7 @@ void GPUTPCGMMerger::DumpMergeCE(std::ostream& out) const
151151
DumpTrackLinks(out, true, " for CE merging");
152152
out << "\nTPC Merger Merge CE\n";
153153
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
154-
const auto& trk = mOutputTracks[i];
154+
const auto& trk = mMergedTracks[i];
155155
if (trk.CCE()) {
156156
out << " Track " << i << ": CCE\n";
157157
}
@@ -167,7 +167,7 @@ void GPUTPCGMMerger::DumpFitPrepare(std::ostream& out) const
167167
}
168168
out << " Clusters\n";
169169
for (uint32_t j = 0; j < mMemory->nMergedTracks; j++) {
170-
const auto& trk = mOutputTracks[j];
170+
const auto& trk = mMergedTracks[j];
171171
out << " Track " << j << ": ";
172172
for (uint32_t i = trk.FirstClusterRef(); i < trk.FirstClusterRef() + trk.NClusters(); i++) {
173173
out << j << "/" << (i - trk.FirstClusterRef()) << ": " << mClusters[i].num << "/" << (int32_t)mClusters[i].state << ", ";
@@ -196,14 +196,14 @@ void GPUTPCGMMerger::DumpRefit(std::ostream& out) const
196196
out << std::setprecision(2);
197197
out << "\nTPC Merger Refit\n";
198198
for (uint32_t i = 0; i < mMemory->nMergedTracks; i++) {
199-
const auto& trk = mOutputTracks[i];
199+
const auto& trk = mMergedTracks[i];
200200
if (trk.NClusters() == 0) {
201201
continue;
202202
}
203203
const auto& p = trk.GetParam();
204204
const auto& po = trk.OuterParam();
205205
out << " Track " << i << ": OK " << trk.OK() << " 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() << " / " << trk.NClustersFitted() << " Cov " << p.GetErr2Y() << "/" << p.GetErr2Z()
206-
<< " dEdx " << (trk.OK() && Param().dodEdxEnabled ? mOutputTracksdEdx[i].dEdxTotTPC : -1.f) << "/" << (trk.OK() && Param().dodEdxEnabled ? mOutputTracksdEdx[i].dEdxMaxTPC : -1.f)
206+
<< " dEdx " << (trk.OK() && Param().dodEdxEnabled ? mMergedTracksdEdx[i].dEdxTotTPC : -1.f) << "/" << (trk.OK() && Param().dodEdxEnabled ? mMergedTracksdEdx[i].dEdxMaxTPC : -1.f)
207207
<< " Outer " << po.P[0] << "/" << po.P[1] << "/" << po.P[2] << "/" << po.P[3] << "/" << po.P[4]
208208
<< " NFitted " << trk.NClustersFitted() << " legs " << (int)trk.Legs() << " flags " << (int)trk.Flags() << "\n";
209209
}
@@ -217,7 +217,7 @@ void GPUTPCGMMerger::DumpLoopers(std::ostream& out) const
217217
if (i && i % 100 == 0) {
218218
out << "\n";
219219
}
220-
out << (int)mOutputTracks[i].MergedLooper() << " ";
220+
out << (int)mMergedTracks[i].MergedLooper() << " ";
221221
}
222222
out << "\n";
223223
}
@@ -226,7 +226,7 @@ void GPUTPCGMMerger::DumpFinal(std::ostream& out) const
226226
{
227227
out << "\nTPC Merger Finalized\n";
228228
for (uint32_t j = 0; j < mMemory->nMergedTracks; j++) {
229-
const auto& trk = mOutputTracks[j];
229+
const auto& trk = mMergedTracks[j];
230230
if (trk.NClusters() == 0) {
231231
continue;
232232
}

‎GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ GPUdii() void GPUTPCGMMergerTrackFit::Thread<0>(int32_t nBlocks, int32_t nThread
2424
const int32_t iEnd = mode == -1 ? merger.Memory()->nRetryRefit : merger.NMergedTracks();
2525
GPUCA_TBB_KERNEL_LOOP(merger.GetRec(), int32_t, ii, iEnd, {
2626
const int32_t i = mode == -1 ? merger.RetryRefitIds()[ii] : mode ? merger.TrackOrderProcess()[ii] : ii;
27-
GPUTPCGMTrackParam::RefitTrack(merger.OutputTracks()[i], i, &merger, mode == -1);
27+
GPUTPCGMTrackParam::RefitTrack(merger.MergedTracks()[i], i, &merger, mode == -1);
2828
});
2929
}
3030

0 commit comments

Comments
 (0)