From 3467dde6c7ab93466067f2a0de34921f96b59b3b Mon Sep 17 00:00:00 2001 From: David Rohr Date: Mon, 21 Jul 2025 15:21:28 +0200 Subject: [PATCH 01/15] GPU Display: Block until display started and print meaningful info messages to the console --- GPU/GPUTracking/Global/GPUChainTracking.cxx | 7 +++- .../Interface/GPUO2InterfaceDisplay.cxx | 6 ++- GPU/GPUTracking/display/GPUDisplay.cxx | 42 ++++++++++++++----- GPU/GPUTracking/display/GPUDisplay.h | 27 +++++++----- GPU/GPUTracking/display/GPUDisplayInterface.h | 3 +- .../display/frontend/GPUDisplayKeys.cxx | 2 +- .../display/render/GPUDisplayDraw.cxx | 2 +- 7 files changed, 61 insertions(+), 28 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx index 9d2578731a30c..79e9ce6cef766 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.cxx +++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx @@ -824,6 +824,7 @@ int32_t GPUChainTracking::RunChainFinalize() if (GetProcessingSettings().eventDisplay) { if (!mDisplayRunning) { + GPUInfo("Starting Event Display..."); if (mEventDisplay->StartDisplay()) { return (1); } @@ -832,6 +833,8 @@ int32_t GPUChainTracking::RunChainFinalize() mEventDisplay->ShowNextEvent(); } + mEventDisplay->WaitTillEventShown(); + if (GetProcessingSettings().eventDisplay->EnableSendKey()) { while (kbhit()) { getch(); @@ -863,9 +866,9 @@ int32_t GPUChainTracking::RunChainFinalize() return (2); } GetProcessingSettings().eventDisplay->setDisplayControl(0); - GPUInfo("Loading next event"); + GPUInfo("Loading next event..."); - mEventDisplay->WaitForNextEvent(); + mEventDisplay->BlockTillNextEvent(); } return 0; diff --git a/GPU/GPUTracking/Interface/GPUO2InterfaceDisplay.cxx b/GPU/GPUTracking/Interface/GPUO2InterfaceDisplay.cxx index 60d5eaf9ae162..ad740200a253a 100644 --- a/GPU/GPUTracking/Interface/GPUO2InterfaceDisplay.cxx +++ b/GPU/GPUTracking/Interface/GPUO2InterfaceDisplay.cxx @@ -46,7 +46,8 @@ int32_t GPUO2InterfaceDisplay::startDisplay() if (retVal) { return retVal; } - mDisplay->WaitForNextEvent(); + mDisplay->WaitTillEventShown(); + mDisplay->BlockTillNextEvent(); return 0; } @@ -59,6 +60,7 @@ int32_t GPUO2InterfaceDisplay::show(const GPUTrackingInOutPointers* ptrs) ptrs = tmpPtr.get(); } mDisplay->ShowNextEvent(ptrs); + mDisplay->WaitTillEventShown(); do { usleep(10000); } while (mFrontend->getDisplayControl() == 0); @@ -66,7 +68,7 @@ int32_t GPUO2InterfaceDisplay::show(const GPUTrackingInOutPointers* ptrs) return 1; } mFrontend->setDisplayControl(0); - mDisplay->WaitForNextEvent(); + mDisplay->BlockTillNextEvent(); return 0; } diff --git a/GPU/GPUTracking/display/GPUDisplay.cxx b/GPU/GPUTracking/display/GPUDisplay.cxx index 136b1947f60ee..163e4c0981bc2 100644 --- a/GPU/GPUTracking/display/GPUDisplay.cxx +++ b/GPU/GPUTracking/display/GPUDisplay.cxx @@ -21,6 +21,8 @@ #include #include #include +#include +#include #ifndef _WIN32 #include "../utils/linux_helpers.h" @@ -143,7 +145,7 @@ void GPUDisplay::ResizeScene(int32_t width, int32_t height, bool init) mBackend->resizeScene(width, height); if (init) { - mResetScene = 1; + mResetScene = true; mViewMatrix = MY_HMM_IDENTITY; mModelMatrix = MY_HMM_IDENTITY; } @@ -220,6 +222,14 @@ int32_t GPUDisplay::DrawGLScene() GPUError("Runtime error %s during display", e.what()); retVal = 1; } + + if (mLoadAndShowEvent) { + { + std::lock_guard lock(mMutexLoadAndShowEvent); + mLoadAndShowEvent = false; + } + mCVLoadAndShowEvent.notify_one(); + } mSemLockDisplay.Unlock(); return retVal; @@ -266,9 +276,9 @@ void GPUDisplay::DrawGLScene_cameraAndAnimation(float animateTime, float& mixSla mCfgR.camLookOrigin = mCfgR.camYUp = false; mAngleRollOrigin = -1e9f; mCfgR.fov = 45.f; - mUpdateDrawCommands = 1; + mUpdateDrawCommands = true; - mResetScene = 0; + mResetScene = false; } else { float moveZ = scalefactor * ((float)mMouseWheelTmp / 150 + (float)(mFrontend->mKeys[(uint8_t)'W'] - mFrontend->mKeys[(uint8_t)'S']) * (!mFrontend->mKeys[mFrontend->KEY_SHIFT]) * 0.2f * mFPSScale); float moveY = scalefactor * ((float)(mFrontend->mKeys[mFrontend->KEY_PAGEDOWN] - mFrontend->mKeys[mFrontend->KEY_PAGEUP]) * 0.2f * mFPSScale); @@ -386,7 +396,7 @@ void GPUDisplay::DrawGLScene_cameraAndAnimation(float animateTime, float& mixSla } if (deltaLine) { SetInfo("%s line width: %f", deltaLine > 0 ? "Increasing" : "Decreasing", mCfgL.lineWidth); - mUpdateDrawCommands = 1; + mUpdateDrawCommands = true; } minSize *= 2; int32_t deltaPoint = mFrontend->mKeys[(uint8_t)'+'] * (!mFrontend->mKeysShift[(uint8_t)'+']) - mFrontend->mKeys[(uint8_t)'-'] * (!mFrontend->mKeysShift[(uint8_t)'-']); @@ -396,7 +406,7 @@ void GPUDisplay::DrawGLScene_cameraAndAnimation(float animateTime, float& mixSla } if (deltaPoint) { SetInfo("%s point size: %f", deltaPoint > 0 ? "Increasing" : "Decreasing", mCfgL.pointSize); - mUpdateDrawCommands = 1; + mUpdateDrawCommands = true; } } @@ -616,7 +626,7 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) disableUnsupportedOptions(); } if (mUpdateEventData || mUpdateVertexLists) { - mUpdateDrawCommands = 1; + mUpdateDrawCommands = true; } if (animateTime < 0 && (mUpdateEventData || mResetScene) && mIOPtrs) { @@ -625,8 +635,8 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) mTimerFPS.ResetStart(); mFramesDoneFPS = 0; mFPSScaleadjust = 0; - mUpdateVertexLists = 1; - mUpdateEventData = 0; + mUpdateVertexLists = true; + mUpdateEventData = false; } hmm_mat4 nextViewMatrix = MY_HMM_IDENTITY; @@ -658,7 +668,7 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) mBackend->drawField(); } - mUpdateDrawCommands = mUpdateRenderPipeline = 0; + mUpdateDrawCommands = mUpdateRenderPipeline = false; mBackend->finishDraw(doScreenshot, renderToMixBuffer, mixSlaveImage); if (animateTime < 0) { @@ -708,15 +718,25 @@ void GPUDisplay::ShowNextEvent(const GPUTrackingInOutPointers* ptrs) if (mMaxClusterZ <= 0) { mResetScene = true; } - mSemLockDisplay.Unlock(); mFrontend->mNeedUpdate = 1; mUpdateEventData = true; + mLoadAndShowEvent = true; + mSemLockDisplay.Unlock(); } -void GPUDisplay::WaitForNextEvent() { mSemLockDisplay.Lock(); } +void GPUDisplay::BlockTillNextEvent() { mSemLockDisplay.Lock(); } + +void GPUDisplay::WaitTillEventShown() +{ + std::unique_lock lock(mMutexLoadAndShowEvent); + while (mLoadAndShowEvent) { + mCVLoadAndShowEvent.wait(lock); + } +} int32_t GPUDisplay::StartDisplay() { + mLoadAndShowEvent = true; if (mFrontend->StartDisplay()) { return (1); } diff --git a/GPU/GPUTracking/display/GPUDisplay.h b/GPU/GPUTracking/display/GPUDisplay.h index 06977c26e0b63..b59e3c52e9bd3 100644 --- a/GPU/GPUTracking/display/GPUDisplay.h +++ b/GPU/GPUTracking/display/GPUDisplay.h @@ -20,13 +20,15 @@ #include "GPUDisplayInterface.h" #include "GPUSettings.h" -#include "../utils/vecpod.h" -#include "../utils/qsem.h" - #include +#include +#include + #include "HandMadeMath.h" #include "utils/timer.h" +#include "utils/vecpod.h" +#include "utils/qsem.h" namespace o2::gpu { @@ -44,7 +46,8 @@ class GPUDisplay : public GPUDisplayInterface int32_t StartDisplay() override; void ShowNextEvent(const GPUTrackingInOutPointers* ptrs = nullptr) override; - void WaitForNextEvent() override; + void BlockTillNextEvent() override; + void WaitTillEventShown() override; void SetCollisionFirstCluster(uint32_t collision, int32_t sector, int32_t cluster) override; void UpdateCalib(const GPUCalibObjectsConst* calib) override { mCalib = calib; } void UpdateParam(const GPUParam* param) override { mParam = param; } @@ -221,7 +224,10 @@ class GPUDisplay : public GPUDisplayInterface GPUSettingsDisplayRenderer mCfgR; const GPUSettingsProcessing& mProcessingSettings; GPUQA* mQA; + qSem mSemLockDisplay; + std::mutex mMutexLoadAndShowEvent; + std::condition_variable mCVLoadAndShowEvent; bool mDrawTextInCompatMode = false; int32_t mDrawTextFontSize = 0; @@ -272,13 +278,14 @@ class GPUDisplay : public GPUDisplayInterface vecpod mTRDTrackIds; vecpod mITSStandaloneTracks; std::vector mTrackFilter; - bool mUpdateTrackFilter = false; - int32_t mUpdateVertexLists = 1; - int32_t mUpdateEventData = 0; - int32_t mUpdateDrawCommands = 1; - int32_t mUpdateRenderPipeline = 0; - volatile int32_t mResetScene = 0; + volatile bool mUpdateTrackFilter = false; + volatile bool mUpdateVertexLists = true; + volatile bool mUpdateEventData = false; + volatile bool mUpdateDrawCommands = true; + volatile bool mUpdateRenderPipeline = false; + volatile bool mResetScene = false; + volatile bool mLoadAndShowEvent = false; int32_t mAnimate = 0; HighResTimer mAnimationTimer; diff --git a/GPU/GPUTracking/display/GPUDisplayInterface.h b/GPU/GPUTracking/display/GPUDisplayInterface.h index 574a8cffc71f0..7caceb1699da6 100644 --- a/GPU/GPUTracking/display/GPUDisplayInterface.h +++ b/GPU/GPUTracking/display/GPUDisplayInterface.h @@ -40,7 +40,8 @@ class GPUDisplayInterface virtual ~GPUDisplayInterface(); virtual int32_t StartDisplay() = 0; virtual void ShowNextEvent(const GPUTrackingInOutPointers* ptrs = nullptr) = 0; - virtual void WaitForNextEvent() = 0; + virtual void BlockTillNextEvent() = 0; + virtual void WaitTillEventShown() = 0; virtual void SetCollisionFirstCluster(uint32_t collision, int32_t sector, int32_t cluster) = 0; virtual void UpdateCalib(const GPUCalibObjectsConst* calib) = 0; virtual void UpdateParam(const GPUParam* param) = 0; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx index 54258857a244c..e1e6d9e54df0a 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx @@ -89,7 +89,7 @@ void GPUDisplay::HandleKey(uint8_t key) mFrontend->mDisplayControl = 2; SetInfo("Exiting", 1); } else if (key == 'r') { - mResetScene = 1; + mResetScene = true; SetInfo("View reset", 1); } else if (key == mFrontend->KEY_ALT && mFrontend->mKeysShift[mFrontend->KEY_ALT]) { mCfgR.camLookOrigin ^= 1; diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index cbe385324ebbf..1dae7c133981e 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -941,7 +941,7 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() break; // TODO: Only sector 0 filled for now } - mUpdateVertexLists = 0; + mUpdateVertexLists = false; size_t totalVertizes = 0; for (int32_t i = 0; i < NSECTORS; i++) { totalVertizes += mVertexBuffer[i].size(); From 373dc30bb6a467ecb577df9476ac9a03e2a5f2f9 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 10:01:55 +0200 Subject: [PATCH 02/15] GPU: Use [[maybe_unused]] to silence compiler warnings --- .../Global/GPUChainTrackingClusterizer.cxx | 3 +-- .../Global/GPUChainTrackingDebugAndProfiling.cxx | 3 +-- GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx | 3 +-- .../SectorTracker/GPUTPCTrackletConstructor.cxx | 3 +-- GPU/GPUTracking/display/render/GPUDisplayDraw.cxx | 4 +--- GPU/GPUTracking/qa/GPUQA.cxx | 6 ++---- GPU/GPUTracking/qa/GPUQAHelper.h | 4 ---- GPU/GPUTracking/utils/qconfig.h | 15 +++++---------- 8 files changed, 12 insertions(+), 29 deletions(-) diff --git a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx index 13455efe6cb47..02e5a51fdbaa2 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx @@ -592,7 +592,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) return ForwardTPCDigits(); } #ifdef GPUCA_TPC_GEOMETRY_O2 - int32_t tpcTimeBinCut = mUpdateNewCalibObjects && mNewCalibValues->newTPCTimeBinCut ? mNewCalibValues->tpcTimeBinCut : param().tpcCutTimeBin; + [[maybe_unused]] int32_t tpcTimeBinCut = mUpdateNewCalibObjects && mNewCalibValues->newTPCTimeBinCut ? mNewCalibValues->tpcTimeBinCut : param().tpcCutTimeBin; mRec->PushNonPersistentMemory(qStr2Tag("TPCCLUST")); const auto& threadContext = GetThreadContext(); const bool doGPU = GetRecoStepsGPU() & RecoStep::TPCClusterFinding; @@ -861,7 +861,6 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput) int32_t firstHBF = (mIOPtrs.settingsTF && mIOPtrs.settingsTF->hasTfStartOrbit) ? mIOPtrs.settingsTF->tfStartOrbit : ((mIOPtrs.tpcZS->sector[iSector].count[0] && mIOPtrs.tpcZS->sector[iSector].nZSPtr[0][0]) ? o2::raw::RDHUtils::getHeartBeatOrbit(*(const o2::header::RAWDataHeader*)mIOPtrs.tpcZS->sector[iSector].zsPtr[0][0]) : 0); uint32_t nBlocks = doGPU ? clusterer.mPmemory->counters.nPagesSubsector : GPUTrackingInOutZS::NENDPOINTS; - (void)tpcTimeBinCut; // TODO: To be used in decoding kernels switch (mCFContext->zsVersion) { default: GPUFatal("Data with invalid TPC ZS mode (%d) received", mCFContext->zsVersion); diff --git a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx index 173d2fb916239..7d790d8e3913f 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx @@ -70,7 +70,7 @@ int32_t GPUChainTracking::DoProfile() fwrite(&bmpFH, 1, sizeof(bmpFH), fp2); fwrite(&bmpIH, 1, sizeof(bmpIH), fp2); - int32_t nEmptySync = 0; + [[maybe_unused]] int32_t nEmptySync = 0; for (uint32_t i = 0; i < bmpheight * ConstructorBlockCount() * ConstructorThreadCount(); i += ConstructorBlockCount() * ConstructorThreadCount()) { int32_t fEmpty = 1; for (uint32_t j = 0; j < ConstructorBlockCount() * ConstructorThreadCount(); j++) { @@ -103,7 +103,6 @@ int32_t GPUChainTracking::DoProfile() } else { nEmptySync = 0; } - (void)nEmptySync; // if (nEmptySync == GPUCA_SCHED_ROW_STEP + 2) break; } diff --git a/GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx b/GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx index 5ca20a39d0462..eeabab6ed395f 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingRefit.cxx @@ -28,8 +28,7 @@ int32_t GPUChainTracking::RunRefit() GPUTrackingRefitProcessor& Refit = processors()->trackingRefit; GPUTrackingRefitProcessor& RefitShadow = doGPU ? processorsShadow()->trackingRefit : Refit; - const auto& threadContext = GetThreadContext(); - (void)threadContext; + [[maybe_unused]] const auto& threadContext = GetThreadContext(); SetupGPUProcessor(&Refit, false); RefitShadow.SetPtrsFromGPUConstantMem(processorsShadow(), doGPU ? &processorsDevice()->param : nullptr); RefitShadow.SetPropagator(doGPU ? processorsShadow()->calibObjects.o2Propagator : GetO2Propagator()); diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx index 2660f6d8cbf44..0b22bfa57c89e 100644 --- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx +++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx @@ -274,7 +274,7 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, } CADEBUG(printf("%14s: SEA PROP ROW %3d X %8.3f -", "", iRow, tParam.X()); for (int32_t i = 0; i < 5; i++) { printf(" %8.3f", tParam.Par()[i]); } printf(" -"); for (int32_t i = 0; i < 15; i++) { printf(" %8.3f", tParam.Cov()[i]); } printf("\n")); - bool found = false; + [[maybe_unused]] bool found = false; float yUncorrected = tParam.GetY(), zUncorrected = tParam.GetZ(); do { if (row.NHits() < 1) { @@ -373,7 +373,6 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/, r.mFirstRow = iRow; } } while (false); - (void)found; if (!found && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer) { uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(tracker.ISector(), iRow, yUncorrected)); if (pad < GPUTPCGeometry::NPads(iRow) && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(tracker.ISector(), iRow, pad)) { diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 1dae7c133981e..4910179eae475 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -129,7 +129,7 @@ GPUDisplay::vboList GPUDisplay::DrawClusters(int32_t iSector, int32_t select, ui if (mOverlayTFClusters.size() > 0 || iCol == 0 || mNCollissions) { const int32_t firstCluster = (mOverlayTFClusters.size() > 1 && iCol > 0) ? mOverlayTFClusters[iCol - 1][iSector] : 0; const int32_t lastCluster = (mOverlayTFClusters.size() > 1 && iCol + 1 < mOverlayTFClusters.size()) ? mOverlayTFClusters[iCol][iSector] : (mParam->par.earlyTpcTransform ? mIOPtrs->nClusterData[iSector] : mIOPtrs->clustersNative ? mIOPtrs->clustersNative->nClustersSector[iSector] : 0); - const bool checkClusterCollision = mQA && mNCollissions && mOverlayTFClusters.size() == 0 && mIOPtrs->clustersNative && mIOPtrs->clustersNative->clustersMCTruth; + [[maybe_unused]] const bool checkClusterCollision = mQA && mNCollissions && mOverlayTFClusters.size() == 0 && mIOPtrs->clustersNative && mIOPtrs->clustersNative->clustersMCTruth; for (int32_t cidInSector = firstCluster; cidInSector < lastCluster; cidInSector++) { const int32_t cid = GET_CID(iSector, cidInSector); #ifdef GPUCA_TPC_GEOMETRY_O2 @@ -139,8 +139,6 @@ GPUDisplay::vboList GPUDisplay::DrawClusters(int32_t iSector, int32_t select, ui continue; } } -#else - (void)checkClusterCollision; #endif if (mCfgH.hideUnmatchedClusters && mQA && mQA->SuppressHit(cid)) { continue; diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx index 2e9c3a1870385..ceab9ef31fee3 100644 --- a/GPU/GPUTracking/qa/GPUQA.cxx +++ b/GPU/GPUTracking/qa/GPUQA.cxx @@ -99,7 +99,7 @@ using namespace o2::gpu; bool unattached = attach == 0; \ float qpt = 0; \ bool lowPt = false; \ - bool mev200 = false; \ + [[maybe_unused]] bool mev200 = false; \ bool mergedLooper = false; \ int32_t id = attach & gputpcgmmergertypes::attachTrackMask; \ if (!unattached) { \ @@ -126,7 +126,6 @@ using namespace o2::gpu; #define CHECK_CLUSTER_STATE_NOCOUNT() \ CHECK_CLUSTER_STATE_INIT() \ - (void)mev200; /* silence unused variable warning*/ \ if (!lowPt && !mergedLooper) { \ GPUTPCClusterRejection::GetProtectionStatus(attach, physics, protect); \ } @@ -1953,8 +1952,7 @@ int32_t GPUQA::DrawQAHistograms(TObjArray* qcout) std::vector colorNums(COLORCOUNT); if (!qcout) { - static int32_t initColorsInitialized = initColors(); - (void)initColorsInitialized; + [[maybe_unused]] static int32_t initColorsInitialized = initColors(); } for (int32_t i = 0; i < COLORCOUNT; i++) { colorNums[i] = qcout ? defaultColorNums[i] : mColors[i]->GetNumber(); diff --git a/GPU/GPUTracking/qa/GPUQAHelper.h b/GPU/GPUTracking/qa/GPUQAHelper.h index a830562119467..e9d98f3e4e305 100644 --- a/GPU/GPUTracking/qa/GPUQAHelper.h +++ b/GPU/GPUTracking/qa/GPUQAHelper.h @@ -105,10 +105,6 @@ class GPUTPCTrkLbl *labelWeight = bestLabel.fWeight; *totalWeight = mTotalWeight; *maxCount = bestLabelCount; - } else { - (void)labelWeight; - (void)totalWeight; - (void)maxCount; } U retVal = bestLabel; if (bestLabelCount < (1.f - mTrackMCMaxFake) * mNCl) { diff --git a/GPU/GPUTracking/utils/qconfig.h b/GPU/GPUTracking/utils/qconfig.h index bc755e583c3b7..a809cc69be501 100644 --- a/GPU/GPUTracking/utils/qconfig.h +++ b/GPU/GPUTracking/utils/qconfig.h @@ -97,12 +97,9 @@ enum qConfigRetVal { qcrOK = 0, #define BeginSubConfig(name, instance, parent, preoptname, preoptnameshort, descr, ...) \ { \ - constexpr const char* preopt = preoptname; \ - (void)preopt; \ - constexpr const char preoptshort = preoptnameshort; \ - (void)preoptshort; \ - name& tmp = parent.instance; \ - (void)tmp; \ + [[maybe_unused]] constexpr const char* preopt = preoptname; \ + [[maybe_unused]] constexpr const char preoptshort = preoptnameshort; \ + [[maybe_unused]] name& tmp = parent.instance; \ bool tmpfound = true; \ if (found) { \ } @@ -174,10 +171,8 @@ enum qConfigRetVal { qcrOK = 0, const char* qon_mxcat(qConfig_subconfig_, name) = preoptnameshort == 0 ? (qon_mxstr(name) ": --" preoptname "\n\t\t" descr) : (qon_mxstr(name) ": -" qon_mxstr('a') " (--" preoptname ")\n\t\t" descr); \ (void)qon_mxcat(qConfig_subconfig_, name); \ if (subConfig == nullptr || strcmp(subConfig, followSub == 2 ? qon_mxstr(name) : preoptname) == 0) { \ - constexpr const char* preopt = preoptname; \ - (void)preopt; \ - constexpr const char preoptshort = preoptnameshort; \ - (void)preoptshort; \ + [[maybe_unused]] constexpr const char* preopt = preoptname; \ + [[maybe_unused]] constexpr const char preoptshort = preoptnameshort; \ char argBuffer[2] = {preoptnameshort, 0}; \ printf("\n %s: (--%s%s%s)\n", descr, preoptname, preoptnameshort == 0 ? "" : " or -", argBuffer); #define BeginHiddenConfig(name, instance) { From a027d11ec2b521c1397e520610e065bf1d571d11 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 10:02:15 +0200 Subject: [PATCH 03/15] GPU: Improve timing messages for GPU Display and GPU QA --- GPU/GPUTracking/display/GPUDisplay.cxx | 15 ++++-- GPU/GPUTracking/display/GPUDisplay.h | 2 +- .../display/render/GPUDisplayDraw.cxx | 53 ++++++++++++++++--- .../display/render/GPUDisplayImportEvent.cxx | 21 +++++++- GPU/GPUTracking/qa/GPUQA.cxx | 29 +++++----- GPU/GPUTracking/utils/timer.cxx | 7 +++ GPU/GPUTracking/utils/timer.h | 1 + 7 files changed, 97 insertions(+), 31 deletions(-) diff --git a/GPU/GPUTracking/display/GPUDisplay.cxx b/GPU/GPUTracking/display/GPUDisplay.cxx index 163e4c0981bc2..e7c04a1bfb407 100644 --- a/GPU/GPUTracking/display/GPUDisplay.cxx +++ b/GPU/GPUTracking/display/GPUDisplay.cxx @@ -237,6 +237,7 @@ int32_t GPUDisplay::DrawGLScene() void GPUDisplay::DrawGLScene_cameraAndAnimation(float animateTime, float& mixSlaveImage, hmm_mat4& nextViewMatrix) { + HighResTimer timer(mUpdateVertexLists && mChain->GetProcessingSettings().debugLevel >= 2); int32_t mMouseWheelTmp = mFrontend->mMouseWheel; mFrontend->mMouseWheel = 0; bool lookOrigin = mCfgR.camLookOrigin ^ mFrontend->mKeys[mFrontend->KEY_ALT]; @@ -420,6 +421,9 @@ void GPUDisplay::DrawGLScene_cameraAndAnimation(float animateTime, float& mixSla mFrontend->mMouseDnX = mFrontend->mMouseMvX; mFrontend->mMouseDnY = mFrontend->mMouseMvY; } + if (timer.IsRunning()) { + GPUInfo("Display Time: Camera:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } } void GPUDisplay::DrawGLScene_drawCommands() @@ -618,7 +622,6 @@ void GPUDisplay::DrawGLScene_drawCommands() void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) // negative time = no mixing { - bool showTimer = false; bool doScreenshot = (mRequestScreenshot || mAnimateScreenshot) && animateTime < 0; updateOptions(); @@ -629,8 +632,9 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) mUpdateDrawCommands = true; } + HighResTimer timerDraw; if (animateTime < 0 && (mUpdateEventData || mResetScene) && mIOPtrs) { - showTimer = true; + timerDraw.ResetStart(); DrawGLScene_updateEventData(); mTimerFPS.ResetStart(); mFramesDoneFPS = 0; @@ -646,8 +650,8 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) // Prepare Event if (mUpdateVertexLists && mIOPtrs) { size_t totalVertizes = DrawGLScene_updateVertexList(); - if (showTimer) { - printf("Event visualization time: %'d us (vertices %'ld / %'ld bytes)\n", (int32_t)(mTimerDraw.GetCurrentElapsedTime() * 1000000.), (int64_t)totalVertizes, (int64_t)(totalVertizes * sizeof(mVertexBuffer[0][0]))); + if (timerDraw.IsRunning()) { + printf("Event visualization time: %'d us (vertices %'ld / %'ld bytes)\n", (int32_t)(timerDraw.GetCurrentElapsedTime() * 1000000.), (int64_t)totalVertizes, (int64_t)(totalVertizes * sizeof(mVertexBuffer[0][0]))); } } @@ -668,7 +672,8 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) mBackend->drawField(); } - mUpdateDrawCommands = mUpdateRenderPipeline = false; + mUpdateDrawCommands = false; + mUpdateRenderPipeline = false; mBackend->finishDraw(doScreenshot, renderToMixBuffer, mixSlaveImage); if (animateTime < 0) { diff --git a/GPU/GPUTracking/display/GPUDisplay.h b/GPU/GPUTracking/display/GPUDisplay.h index b59e3c52e9bd3..c8deeb2378970 100644 --- a/GPU/GPUTracking/display/GPUDisplay.h +++ b/GPU/GPUTracking/display/GPUDisplay.h @@ -310,7 +310,7 @@ class GPUDisplay : public GPUDisplayInterface float mFPSScale = 1, mFPSScaleadjust = 0; int32_t mFramesDone = 0, mFramesDoneFPS = 0; - HighResTimer mTimerFPS, mTimerDisplay, mTimerDraw; + HighResTimer mTimerFPS; vboList mGlDLLines[NSECTORS][N_LINES_TYPE]; vecpod> mGlDLFinal[NSECTORS]; vboList mGlDLFinalITS; diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 4910179eae475..9d524ce16c89b 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -752,6 +752,7 @@ GPUDisplay::vboList GPUDisplay::DrawGridTRD(int32_t sector) size_t GPUDisplay::DrawGLScene_updateVertexList() { + HighResTimer timer(mChain->GetProcessingSettings().debugLevel >= 2); for (int32_t i = 0; i < NSECTORS; i++) { mVertexBuffer[i].clear(); mVertexBufferStart[i].clear(); @@ -773,6 +774,10 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() mGlDLFinal[iSector].resize(mNCollissions); } } + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Init:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } + int32_t numThreads = getNumThreads(); tbb::task_arena(numThreads).execute([&] { if (mChain && (mChain->GetRecoSteps() & GPUDataTypes::RecoStep::TPCSectorTracking)) { @@ -782,6 +787,9 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() mGlDLLines[iSector][tINITLINK] = DrawLinks(tracker, tINITLINK, true); tracker.SetPointersDataLinks(mChain->rec()->Res(tracker.MemoryResLinks()).Ptr()); // clang-format off }, tbb::simple_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Links:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } tbb::parallel_for(0, NSECTORS, [&](int32_t iSector) { const GPUTPCTracker& tracker = sectorTracker(iSector); @@ -795,11 +803,17 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() mGlDLGridTRD[iSector] = DrawGridTRD(iSector); } // clang-format off }, tbb::simple_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Seeds:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } tbb::parallel_for(0, NSECTORS, [&](int32_t iSector) { const GPUTPCTracker& tracker = sectorTracker(iSector); mGlDLLines[iSector][tEXTRAPOLATEDTRACK] = DrawTracks(tracker, 1); // clang-format off }, tbb::simple_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Sector Tracks:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } } tbb::parallel_for(0, numThreads, [&](int32_t iThread) { mThreadTracks[iThread].resize(mNCollissions); @@ -872,6 +886,9 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() mThreadTracks[GPUReconstruction::getHostThreadIndex()][col][sector][1].emplace_back(i); }); } + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Sort merged tracks:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } GPUTPCGMPropagator prop; prop.SetMaxSinPhi(.999); @@ -900,6 +917,9 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() } } // clang-format off }, tbb::simple_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Merged Tracks:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } tbb::parallel_for(0, NSECTORS, [&](int32_t iSector) { for (int32_t i = 0; i < N_POINTS_TYPE_TPC; i++) { @@ -908,36 +928,49 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() } } // clang-format off }, tbb::simple_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Clusters:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } + }); // End omp parallel mGlDLFinalITS = DrawFinalITS(); - for (int32_t iSector = 0; iSector < NSECTORS; iSector++) { - for (int32_t i = N_POINTS_TYPE_TPC; i < N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD; i++) { + for (int32_t i = N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD + N_POINTS_TYPE_TOF; i < N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD + N_POINTS_TYPE_TOF + N_POINTS_TYPE_ITS; i++) { for (int32_t iCol = 0; iCol < mNCollissions; iCol++) { - mGlDLPoints[iSector][i][iCol] = DrawSpacePointsTRD(iSector, i, iCol); + mGlDLPoints[iSector][i][iCol] = DrawSpacePointsITS(iSector, i, iCol); } } + break; // TODO: Only sector 0 filled for now + } + + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex ITS:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } for (int32_t iSector = 0; iSector < NSECTORS; iSector++) { - for (int32_t i = N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD; i < N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD + N_POINTS_TYPE_TOF; i++) { + for (int32_t i = N_POINTS_TYPE_TPC; i < N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD; i++) { for (int32_t iCol = 0; iCol < mNCollissions; iCol++) { - mGlDLPoints[iSector][i][iCol] = DrawSpacePointsTOF(iSector, i, iCol); + mGlDLPoints[iSector][i][iCol] = DrawSpacePointsTRD(iSector, i, iCol); } } - break; // TODO: Only sector 0 filled for now + } + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex TRD:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } for (int32_t iSector = 0; iSector < NSECTORS; iSector++) { - for (int32_t i = N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD + N_POINTS_TYPE_TOF; i < N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD + N_POINTS_TYPE_TOF + N_POINTS_TYPE_ITS; i++) { + for (int32_t i = N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD; i < N_POINTS_TYPE_TPC + N_POINTS_TYPE_TRD + N_POINTS_TYPE_TOF; i++) { for (int32_t iCol = 0; iCol < mNCollissions; iCol++) { - mGlDLPoints[iSector][i][iCol] = DrawSpacePointsITS(iSector, i, iCol); + mGlDLPoints[iSector][i][iCol] = DrawSpacePointsTOF(iSector, i, iCol); } } break; // TODO: Only sector 0 filled for now } + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex TOF:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } mUpdateVertexLists = false; size_t totalVertizes = 0; @@ -965,5 +998,9 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() for (int32_t i = 0; i < (mUseMultiVBO ? GPUCA_NSECTORS : 1); i++) { mVertexBuffer[i].clear(); } + if (timer.IsRunning()) { + GPUInfo("Display Time: Vertex Final:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } + return totalVertizes; } diff --git a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx index ab4c0abd7b60e..df3b385c14fe5 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx @@ -23,6 +23,7 @@ #include "GPUTPCConvertImpl.h" #include "GPUTRDGeometry.h" #include "GPUTRDTrackletWord.h" +#include "GPUChainTracking.h" #include "GPUParam.inc" #include "DataFormatsTOF/Cluster.h" @@ -40,7 +41,7 @@ using namespace o2::gpu; void GPUDisplay::DrawGLScene_updateEventData() { - mTimerDraw.ResetStart(); + HighResTimer timer(mChain->GetProcessingSettings().debugLevel >= 2); if (mIOPtrs->clustersNative) { mCurrentClusters = mIOPtrs->clustersNative->nClustersTotal; } else { @@ -110,6 +111,9 @@ void GPUDisplay::DrawGLScene_updateEventData() } } } + if (timer.IsRunning()) { + GPUInfo("Display Time: Init:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } if (mCfgH.trackFilter) { uint32_t nTracks = mConfig.showTPCTracksFromO2Format ? mIOPtrs->nOutputTracksTPCO2 : mIOPtrs->nMergedTracks; @@ -128,6 +132,9 @@ void GPUDisplay::DrawGLScene_updateEventData() } } mUpdateTrackFilter = false; + if (timer.IsRunning()) { + GPUInfo("Display Time: Track Filter:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } mMaxClusterZ = tbb::parallel_reduce(tbb::blocked_range(0, NSECTORS, 1), float(0.f), [&](const tbb::blocked_range& r, float maxClusterZ) { for (int32_t iSector = r.begin(); iSector < r.end(); iSector++) { @@ -174,6 +181,9 @@ void GPUDisplay::DrawGLScene_updateEventData() } return maxClusterZ; // clang-format off }, [](const float a, const float b) { return std::max(a, b); }, tbb::simple_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Load TPC:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } mMaxClusterZ = tbb::parallel_reduce(tbb::blocked_range(0, mCurrentSpacePointsTRD, 32), float(mMaxClusterZ), [&](const tbb::blocked_range& r, float maxClusterZ) { int32_t trdTriggerRecord = -1; @@ -209,6 +219,9 @@ void GPUDisplay::DrawGLScene_updateEventData() } return maxClusterZ; // clang-format off }, [](const float a, const float b) { return std::max(a, b); }, tbb::static_partitioner()); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Load TRD:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } mMaxClusterZ = tbb::parallel_reduce(tbb::blocked_range(0, mCurrentClustersTOF, 32), float(mMaxClusterZ), [&](const tbb::blocked_range& r, float maxClusterZ) { for (int32_t i = r.begin(); i < r.end(); i++) { @@ -230,6 +243,9 @@ void GPUDisplay::DrawGLScene_updateEventData() } return maxClusterZ; // clang-format off }, [](const float a, const float b) { return std::max(a, b); }); // clang-format on + if (timer.IsRunning()) { + GPUInfo("Display Time: Load TOF:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } if (mCurrentClustersITS) { float itsROFhalfLen = 0; @@ -270,4 +286,7 @@ void GPUDisplay::DrawGLScene_updateEventData() } } } + if (timer.IsRunning()) { + GPUInfo("Display Time: Load ITS:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); + } } diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx index ceab9ef31fee3..673f63793939d 100644 --- a/GPU/GPUTracking/qa/GPUQA.cxx +++ b/GPU/GPUTracking/qa/GPUQA.cxx @@ -661,10 +661,9 @@ void GPUQA::InitO2MCData(GPUTrackingInOutPointers* updateIOPtr) { #ifdef GPUCA_O2_LIB if (!mO2MCDataLoaded) { - HighResTimer timer; + HighResTimer timer(mTracking && mTracking->GetProcessingSettings().debugLevel); if (mTracking && mTracking->GetProcessingSettings().debugLevel) { GPUInfo("Start reading O2 Track MC information"); - timer.Start(); } static constexpr float PRIM_MAX_T = 0.01f; @@ -756,7 +755,7 @@ void GPUQA::InitO2MCData(GPUTrackingInOutPointers* updateIOPtr) } } } - if (mTracking && mTracking->GetProcessingSettings().debugLevel) { + if (timer.IsRunning()) { GPUInfo("Finished reading O2 Track MC information (%f seconds)", timer.GetCurrentElapsedTime()); } mO2MCDataLoaded = true; @@ -902,7 +901,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx mClusterParam.resize(GetNMCLabels()); memset(mClusterParam.data(), 0, mClusterParam.size() * sizeof(mClusterParam[0])); } - HighResTimer timer; + HighResTimer timer(QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 2)); mNEvents++; if (mConfig.writeMCLabels) { @@ -914,9 +913,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx bool mcAvail = mcPresent() || tracksExtMC; - if (mcAvail) { - // Assign Track MC Labels - timer.Start(); + if (mcAvail) { // Assign Track MC Labels if (tracksExternal) { #ifdef GPUCA_O2_LIB for (uint32_t i = 0; i < tracksExternal->size(); i++) { @@ -967,7 +964,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } }); } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Assign Track Labels:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } @@ -1107,7 +1104,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Cluster attach status:\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } @@ -1134,7 +1131,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Compute cluster label weights:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } @@ -1158,7 +1155,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } // clang-format off }, tbb::simple_partitioner()); // clang-format on - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Compute track mc parameters:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } @@ -1248,7 +1245,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Fill efficiency histograms:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } } @@ -1421,7 +1418,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Fill resolution histograms:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } } @@ -1648,7 +1645,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Fill cluster histograms:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } } @@ -1737,7 +1734,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx clusterAttachCounts.clear(); } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Fill track statistics:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } } @@ -1810,7 +1807,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx mClusterCounts = counts_t(); } - if (QA_TIMING || (mTracking && mTracking->GetProcessingSettings().debugLevel >= 3)) { + if (timer.IsRunning()) { GPUInfo("QA Time: Cluster Counts:\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } diff --git a/GPU/GPUTracking/utils/timer.cxx b/GPU/GPUTracking/utils/timer.cxx index f3b108fc6f159..df3790ad9ccbf 100644 --- a/GPU/GPUTracking/utils/timer.cxx +++ b/GPU/GPUTracking/utils/timer.cxx @@ -23,6 +23,13 @@ #include #endif +HighResTimer::HighResTimer(bool start) +{ + if (start) { + ResetStart(); + } +} + inline double HighResTimer::GetTime() { #ifdef _WIN32 diff --git a/GPU/GPUTracking/utils/timer.h b/GPU/GPUTracking/utils/timer.h index 44a01b04747cb..35b1d707b97b0 100644 --- a/GPU/GPUTracking/utils/timer.h +++ b/GPU/GPUTracking/utils/timer.h @@ -21,6 +21,7 @@ class HighResTimer { public: HighResTimer() = default; + HighResTimer(bool start); ~HighResTimer() = default; void Start(); void Stop(); From 01aa0c54b48bbd067f478d2ae4d30356481e75db Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 10:50:58 +0200 Subject: [PATCH 04/15] GPU Display: Speed up drawing clusters with many collisions --- GPU/GPUTracking/display/GPUDisplay.h | 3 +- .../display/render/GPUDisplayDraw.cxx | 135 ++++++++++-------- 2 files changed, 78 insertions(+), 60 deletions(-) diff --git a/GPU/GPUTracking/display/GPUDisplay.h b/GPU/GPUTracking/display/GPUDisplay.h index c8deeb2378970..1cdbf62da2202 100644 --- a/GPU/GPUTracking/display/GPUDisplay.h +++ b/GPU/GPUTracking/display/GPUDisplay.h @@ -193,7 +193,7 @@ class GPUDisplay : public GPUDisplayInterface void SetCollisionColor(int32_t col); void updateConfig(); void drawPointLinestrip(int32_t iSector, int32_t cid, int32_t id, int32_t id_limit = TRACK_TYPE_ID_LIMIT); - vboList DrawClusters(int32_t iSector, int32_t select, uint32_t iCol); + void DrawClusters(int32_t iSector); vboList DrawSpacePointsTRD(int32_t iSector, int32_t select, int32_t iCol); vboList DrawSpacePointsTOF(int32_t iSector, int32_t select, int32_t iCol); vboList DrawSpacePointsITS(int32_t iSector, int32_t select, int32_t iCol); @@ -256,6 +256,7 @@ class GPUDisplay : public GPUDisplayInterface vecpod mVertexBuffer[NSECTORS]; vecpod mVertexBufferStart[NSECTORS]; vecpod mVertexBufferCount[NSECTORS]; + std::vector> mClusterBufferSizeCache[NSECTORS]; std::unique_ptr mGlobalPosPtr; std::unique_ptr mGlobalPosPtrTRD; diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 9d524ce16c89b..baf57175ecea8 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -122,70 +122,91 @@ GPUDisplay::vboList GPUDisplay::DrawSpacePointsITS(int32_t iSector, int32_t sele return (vboList(startCount, mVertexBufferStart[iSector].size() - startCount, iSector)); } -GPUDisplay::vboList GPUDisplay::DrawClusters(int32_t iSector, int32_t select, uint32_t iCol) +void GPUDisplay::DrawClusters(int32_t iSector) { - size_t startCount = mVertexBufferStart[iSector].size(); - size_t startCountInner = mVertexBuffer[iSector].size(); - if (mOverlayTFClusters.size() > 0 || iCol == 0 || mNCollissions) { - const int32_t firstCluster = (mOverlayTFClusters.size() > 1 && iCol > 0) ? mOverlayTFClusters[iCol - 1][iSector] : 0; - const int32_t lastCluster = (mOverlayTFClusters.size() > 1 && iCol + 1 < mOverlayTFClusters.size()) ? mOverlayTFClusters[iCol][iSector] : (mParam->par.earlyTpcTransform ? mIOPtrs->nClusterData[iSector] : mIOPtrs->clustersNative ? mIOPtrs->clustersNative->nClustersSector[iSector] : 0); - [[maybe_unused]] const bool checkClusterCollision = mQA && mNCollissions && mOverlayTFClusters.size() == 0 && mIOPtrs->clustersNative && mIOPtrs->clustersNative->clustersMCTruth; - for (int32_t cidInSector = firstCluster; cidInSector < lastCluster; cidInSector++) { - const int32_t cid = GET_CID(iSector, cidInSector); + std::vector, N_POINTS_TYPE_TPC>> vertexCache(mNCollissions); + if (mClusterBufferSizeCache[iSector].size() < (uint32_t)mNCollissions) { + mClusterBufferSizeCache[iSector].resize(mNCollissions); + } + for (int32_t iCol = 0; iCol < mNCollissions; iCol++) { + for (int32_t i = 0; i < N_POINTS_TYPE_TPC; i++) { + vertexCache[iCol][i].reserve(mClusterBufferSizeCache[iSector][iCol][i]); + } + } + + uint32_t col = 0; + const int32_t nClustersInSector = mParam->par.earlyTpcTransform ? mIOPtrs->nClusterData[iSector] : (mIOPtrs->clustersNative ? mIOPtrs->clustersNative->nClustersSector[iSector] : 0); + [[maybe_unused]] const bool checkClusterCollision = mQA && mNCollissions && mOverlayTFClusters.size() == 0 && mIOPtrs->clustersNative && mIOPtrs->clustersNative->clustersMCTruth; + for (int32_t cidInSector = 0; cidInSector < nClustersInSector; cidInSector++) { + const int32_t cid = GET_CID(iSector, cidInSector); #ifdef GPUCA_TPC_GEOMETRY_O2 - if (checkClusterCollision) { - const auto& labels = mIOPtrs->clustersNative->clustersMCTruth->getLabels(cid); - if (labels.size() ? (iCol != mQA->GetMCLabelCol(labels[0])) : (iCol != 0)) { - continue; - } - } + if (checkClusterCollision) { + const auto& labels = mIOPtrs->clustersNative->clustersMCTruth->getLabels(cid); + col = labels.size() ? mQA->GetMCLabelCol(labels[0]) : 0; + } else #endif - if (mCfgH.hideUnmatchedClusters && mQA && mQA->SuppressHit(cid)) { - continue; + if (mOverlayTFClusters.size()) { + while (col < mOverlayTFClusters.size() && cidInSector >= mOverlayTFClusters[col][iSector]) { + col++; } - bool draw = mGlobalPos[cid].w == select; - - if (mCfgH.markAdjacentClusters) { - const int32_t attach = mIOPtrs->mergedTrackHitAttachment[cid]; - if (attach) { - if (mCfgH.markAdjacentClusters >= 32) { - if (mQA && mQA->clusterRemovable(attach, mCfgH.markAdjacentClusters == 33)) { - draw = select == tMARKED; - } - } else if ((mCfgH.markAdjacentClusters & 2) && (attach & gputpcgmmergertypes::attachTube)) { - draw = select == tMARKED; - } else if ((mCfgH.markAdjacentClusters & 1) && (attach & (gputpcgmmergertypes::attachGood | gputpcgmmergertypes::attachTube)) == 0) { - draw = select == tMARKED; - } else if ((mCfgH.markAdjacentClusters & 4) && (attach & gputpcgmmergertypes::attachGoodLeg) == 0) { - draw = select == tMARKED; - } else if ((mCfgH.markAdjacentClusters & 16) && (attach & gputpcgmmergertypes::attachHighIncl)) { - draw = select == tMARKED; - } else if (mCfgH.markAdjacentClusters & 8) { - if (fabsf(mIOPtrs->mergedTracks[attach & gputpcgmmergertypes::attachTrackMask].GetParam().GetQPt()) > 20.f) { - draw = select == tMARKED; - } + } + if (mCfgH.hideUnmatchedClusters && mQA && mQA->SuppressHit(cid)) { + continue; + } + int32_t select = mGlobalPos[cid].w; + + if (mCfgH.markAdjacentClusters) { + const int32_t attach = mIOPtrs->mergedTrackHitAttachment[cid]; + if (attach) { + if (mCfgH.markAdjacentClusters >= 32) { + if (mQA && mQA->clusterRemovable(attach, mCfgH.markAdjacentClusters == 33)) { + select = tMARKED; + } + } else if ((mCfgH.markAdjacentClusters & 2) && (attach & gputpcgmmergertypes::attachTube)) { + select = tMARKED; + } else if ((mCfgH.markAdjacentClusters & 1) && (attach & (gputpcgmmergertypes::attachGood | gputpcgmmergertypes::attachTube)) == 0) { + select = tMARKED; + } else if ((mCfgH.markAdjacentClusters & 4) && (attach & gputpcgmmergertypes::attachGoodLeg) == 0) { + select = tMARKED; + } else if ((mCfgH.markAdjacentClusters & 16) && (attach & gputpcgmmergertypes::attachHighIncl)) { + select = tMARKED; + } else if (mCfgH.markAdjacentClusters & 8) { + if (fabsf(mIOPtrs->mergedTracks[attach & gputpcgmmergertypes::attachTrackMask].GetParam().GetQPt()) > 20.f) { + select = tMARKED; } } - } else if (mCfgH.markClusters) { - int16_t flags; - if (mParam->par.earlyTpcTransform) { - flags = mIOPtrs->clusterData[iSector][cidInSector].flags; - } else { - flags = mIOPtrs->clustersNative->clustersLinear[cid].getFlags(); - } - const bool match = flags & mCfgH.markClusters; - draw = (select == tMARKED) ? (match) : (draw && !match); - } else if (mCfgH.markFakeClusters) { - const bool fake = (mQA->HitAttachStatus(cid)); - draw = (select == tMARKED) ? (fake) : (draw && !fake); } - if (draw) { - mVertexBuffer[iSector].emplace_back(mGlobalPos[cid].x, mGlobalPos[cid].y * mYFactor, mCfgH.projectXY ? 0 : mGlobalPos[cid].z); + } else if (mCfgH.markClusters) { + int16_t flags; + if (mParam->par.earlyTpcTransform) { + flags = mIOPtrs->clusterData[iSector][cidInSector].flags; + } else { + flags = mIOPtrs->clustersNative->clustersLinear[cid].getFlags(); + } + if (flags & mCfgH.markClusters) { + select = tMARKED; + } + } else if (mCfgH.markFakeClusters) { + if (mQA->HitAttachStatus(cid)) { + select = tMARKED; } } + vertexCache[col][select].emplace_back(mGlobalPos[cid].x, mGlobalPos[cid].y * mYFactor, mCfgH.projectXY ? 0 : mGlobalPos[cid].z); + } + + size_t startCountInner = mVertexBuffer[iSector].size(); + mVertexBuffer[iSector].resize(mVertexBuffer[iSector].size() + nClustersInSector); + for (int32_t iCol = 0; iCol < mNCollissions; iCol++) { + for (int32_t i = 0; i < N_POINTS_TYPE_TPC; i++) { + uint32_t count = vertexCache[iCol][i].size(); + mClusterBufferSizeCache[iSector][iCol][i] = std::max(mClusterBufferSizeCache[iSector][iCol][i], count); + memcpy((void*)&mVertexBuffer[iSector][startCountInner], (const void*)vertexCache[iCol][i].data(), count * sizeof(vertexCache[iCol][i][0])); + size_t startCount = mVertexBufferStart[iSector].size(); + insertVertexList(iSector, startCountInner, startCountInner + count); + startCountInner += count; + mGlDLPoints[iSector][i][iCol] = vboList(startCount, mVertexBufferStart[iSector].size() - startCount, iSector); + } } - insertVertexList(iSector, startCountInner, mVertexBuffer[iSector].size()); - return (vboList(startCount, mVertexBufferStart[iSector].size() - startCount, iSector)); } GPUDisplay::vboList GPUDisplay::DrawLinks(const GPUTPCTracker& tracker, int32_t id, bool dodown) @@ -922,11 +943,7 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() } tbb::parallel_for(0, NSECTORS, [&](int32_t iSector) { - for (int32_t i = 0; i < N_POINTS_TYPE_TPC; i++) { - for (int32_t iCol = 0; iCol < mNCollissions; iCol++) { - mGlDLPoints[iSector][i][iCol] = DrawClusters(iSector, i, iCol); - } - } // clang-format off + DrawClusters(iSector); // clang-format off }, tbb::simple_partitioner()); // clang-format on if (timer.IsRunning()) { GPUInfo("Display Time: Vertex Clusters:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); From 038b3f662e4c2b07cfdd73cdbfb9bfc5081dc8fc Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 10:54:41 +0200 Subject: [PATCH 05/15] GPU Display: ResetScene should reset which collision to show --- GPU/GPUTracking/display/GPUDisplay.cxx | 1 + 1 file changed, 1 insertion(+) diff --git a/GPU/GPUTracking/display/GPUDisplay.cxx b/GPU/GPUTracking/display/GPUDisplay.cxx index e7c04a1bfb407..35ebb132398ab 100644 --- a/GPU/GPUTracking/display/GPUDisplay.cxx +++ b/GPU/GPUTracking/display/GPUDisplay.cxx @@ -273,6 +273,7 @@ void GPUDisplay::DrawGLScene_cameraAndAnimation(float animateTime, float& mixSla mCfgL.pointSize = 2.0f; mCfgL.lineWidth = 1.4f; mCfgL.drawSector = -1; + mCfgL.showCollision = -1; mCfgH.xAdd = mCfgH.zAdd = 0; mCfgR.camLookOrigin = mCfgR.camYUp = false; mAngleRollOrigin = -1e9f; From 0d593016d644417b70cad107398b436a1fe8342a Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 11:11:27 +0200 Subject: [PATCH 06/15] GPU Display: Extrapolate tracks only on-demand when first requested --- GPU/GPUTracking/display/GPUDisplay.cxx | 2 +- GPU/GPUTracking/display/GPUDisplay.h | 1 + GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx | 3 +++ GPU/GPUTracking/display/render/GPUDisplayDraw.cxx | 4 ++++ 4 files changed, 9 insertions(+), 1 deletion(-) diff --git a/GPU/GPUTracking/display/GPUDisplay.cxx b/GPU/GPUTracking/display/GPUDisplay.cxx index 35ebb132398ab..7cad25916940a 100644 --- a/GPU/GPUTracking/display/GPUDisplay.cxx +++ b/GPU/GPUTracking/display/GPUDisplay.cxx @@ -633,7 +633,7 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) mUpdateDrawCommands = true; } - HighResTimer timerDraw; + HighResTimer timerDraw(mUpdateVertexLists); if (animateTime < 0 && (mUpdateEventData || mResetScene) && mIOPtrs) { timerDraw.ResetStart(); DrawGLScene_updateEventData(); diff --git a/GPU/GPUTracking/display/GPUDisplay.h b/GPU/GPUTracking/display/GPUDisplay.h index 1cdbf62da2202..7279f2ee87fdb 100644 --- a/GPU/GPUTracking/display/GPUDisplay.h +++ b/GPU/GPUTracking/display/GPUDisplay.h @@ -287,6 +287,7 @@ class GPUDisplay : public GPUDisplayInterface volatile bool mUpdateRenderPipeline = false; volatile bool mResetScene = false; volatile bool mLoadAndShowEvent = false; + bool mTracksArePropagated = false; int32_t mAnimate = 0; HighResTimer mAnimationTimer; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx index e1e6d9e54df0a..6dc09545733fe 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx @@ -493,6 +493,9 @@ void GPUDisplay::HandleKey(uint8_t key) if (memcmp((void*)&oldCfgH, (void*)&mCfgH, sizeof(mCfgH)) != 0) { mUpdateEventData = true; } + if (mCfgL.propagateTracks != 0 && !mTracksArePropagated) { + mUpdateVertexLists = true; + } if (memcmp((void*)&oldCfgL, (void*)&mCfgL, sizeof(mCfgL)) != 0 || memcmp((void*)&oldCfgR, (void*)&mCfgR, sizeof(mCfgR)) != 0) { mUpdateDrawCommands = true; } diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index baf57175ecea8..f50e6590eb16d 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -494,6 +494,9 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp if (!mIOPtrs->clustersNative) { continue; } + if (mCfgL.propagateTracks == 0) { + continue; + } // Propagate track paramters / plot MC tracks for (int32_t iMC = 0; iMC < 2; iMC++) { @@ -989,6 +992,7 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() GPUInfo("Display Time: Vertex TOF:\t\t\t%6.0f us", timer.GetCurrentElapsedTime(true) * 1e6); } + mTracksArePropagated = mCfgL.propagateTracks != 0; mUpdateVertexLists = false; size_t totalVertizes = 0; for (int32_t i = 0; i < NSECTORS; i++) { From 4b137481f3ffc81a4e1e48249f9776428ece2a52 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 11:34:16 +0200 Subject: [PATCH 07/15] TPC: Change some default settings --- GPU/GPUTracking/Definitions/GPUSettingsList.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 51a0add7dbeea..59d48f0fd01b3 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -132,11 +132,11 @@ AddOptionRTC(cfInnerThreshold, uint8_t, 0, "", 0, "Cluster Finder extends cluste AddOptionRTC(cfMinSplitNum, uint8_t, 1, "", 0, "Minimum number of split charges in a cluster for the cluster to be marked as split") AddOptionRTC(cfNoiseSuppressionEpsilon, uint8_t, 10, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression") AddOptionRTC(cfNoiseSuppressionEpsilonRelative, uint8_t, 76, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression, relative as fraction of 255") -AddOptionRTC(cfEdgeTwoPads, uint8_t, 1, "", 0, "Flag clusters with peak on the 2 pads closes to the sector edge as edge cluster") +AddOptionRTC(cfEdgeTwoPads, uint8_t, 0, "", 0, "Flag clusters with peak on the 2 pads closes to the sector edge as edge cluster") AddOptionRTC(nWays, uint8_t, 3, "", 0, "Do N fit passes in final fit of merger") AddOptionRTC(nWaysOuter, int8_t, 0, "", 0, "Store outer param") AddOptionRTC(trackFitRejectMode, int8_t, 5, "", 0, "0: no limit on rejection or missed hits, >0: break after n rejected hits, <0: reject at max -n hits") -AddOptionRTC(rejectIFCLowRadiusCluster, uint8_t, 0, "", 0, "Reject clusters that get the IFC mask error during refit") +AddOptionRTC(rejectIFCLowRadiusCluster, uint8_t, 1, "", 0, "Reject clusters that get the IFC mask error during refit") AddOptionRTC(dEdxTruncLow, uint8_t, 2, "", 0, "Low truncation threshold, fraction of 128") AddOptionRTC(dEdxTruncHigh, uint8_t, 77, "", 0, "High truncation threshold, fraction of 128") AddOptionRTC(extrapolationTracking, int8_t, 1, "", 0, "Enable Extrapolation Tracking (prolong tracks to adjacent sectors to find short segments)") From b7fb1cde380485e07ca0f44819daff8c87ae6e6a Mon Sep 17 00:00:00 2001 From: David Rohr Date: Tue, 22 Jul 2025 11:56:53 +0200 Subject: [PATCH 08/15] GPU TPC: Do looper cluster attachment always in separate kernel --- GPU/GPUTracking/Base/GPUReconstruction.cxx | 3 - .../Definitions/GPUDefParametersDefaults.h | 10 ---- GPU/GPUTracking/Definitions/GPUSettingsList.h | 1 - .../Global/GPUChainTrackingMerger.cxx | 4 +- GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx | 57 +++---------------- GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h | 6 +- .../Standalone/Benchmark/standalone.cxx | 1 - GPU/GPUTracking/kernels.cmake | 1 - 8 files changed, 12 insertions(+), 71 deletions(-) diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index c7b61a976021a..0b8da7027c247 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -271,9 +271,6 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice() #endif mProcessingSettings->overrideClusterizerFragmentLen = TPC_MAX_FRAGMENT_LEN_GPU; param().rec.tpc.nWaysOuter = true; - if (param().rec.tpc.looperInterpolationInExtraPass == -1) { - param().rec.tpc.looperInterpolationInExtraPass = 0; - } if (GetProcessingSettings().createO2Output > 1) { mProcessingSettings->createO2Output = 1; } diff --git a/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h b/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h index 48d00b274dc9c..b1f12034d9c2f 100644 --- a/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h +++ b/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h @@ -80,7 +80,6 @@ #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20 #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1 #define GPUCA_PAR_SORT_BEFORE_FIT 1 - #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1 #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1 #define GPUCA_PAR_DEDX_STORAGE_TYPE uint16_t #define GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE half @@ -143,7 +142,6 @@ #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20 #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1 #define GPUCA_PAR_SORT_BEFORE_FIT 1 - #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1 #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1 #define GPUCA_PAR_DEDX_STORAGE_TYPE uint16_t #define GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE half @@ -206,7 +204,6 @@ #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20 #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1 #define GPUCA_PAR_SORT_BEFORE_FIT 1 - #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1 #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1 #define GPUCA_PAR_DEDX_STORAGE_TYPE uint16_t #define GPUCA_PAR_MERGER_INTERPOLATION_ERROR_TYPE half @@ -261,7 +258,6 @@ #define GPUCA_PAR_TRACKLET_SELECTOR_HITS_REG_SIZE 20 #define GPUCA_PAR_ALTERNATE_BORDER_SORT 1 #define GPUCA_PAR_SORT_BEFORE_FIT 1 - #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 1 #define GPUCA_PAR_NO_ATOMIC_PRECHECK 1 #define GPUCA_PAR_COMP_GATHER_KERNEL 4 #define GPUCA_PAR_COMP_GATHER_MODE 3 @@ -529,9 +525,6 @@ #ifndef GPUCA_PAR_SORT_BEFORE_FIT #define GPUCA_PAR_SORT_BEFORE_FIT 0 #endif - #ifndef GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION - #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 0 - #endif #ifndef GPUCA_PAR_COMP_GATHER_KERNEL #define GPUCA_PAR_COMP_GATHER_KERNEL 0 #endif @@ -566,9 +559,6 @@ #ifndef GPUCA_PAR_SORT_BEFORE_FIT #define GPUCA_PAR_SORT_BEFORE_FIT 0 #endif - #ifndef GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION - #define GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION 0 - #endif #ifndef GPUCA_PAR_COMP_GATHER_KERNEL #define GPUCA_PAR_COMP_GATHER_KERNEL 0 #endif diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 59d48f0fd01b3..96b1ad194e39a 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -154,7 +154,6 @@ AddOptionRTC(mergerInterpolateErrors, uint8_t, 1, "", 0, "Use interpolation inst AddOptionRTC(mergerInterpolateRejectAlsoOnCurrentPosition, uint8_t, 1, "", 0, "When using mergerInterpolateErrors, reject based on chi2 twice computed with interpolated and current track position") AddOptionRTC(mergeCE, uint8_t, 1, "", 0, "Merge tracks accross the central electrode") AddOptionRTC(retryRefit, int8_t, 1, "", 0, "Retry refit with seeding errors and without cluster rejection when fit fails (=2 means retry in same kernel, =1 for separate kernel") -AddOptionRTC(looperInterpolationInExtraPass, int8_t, -1, "", 0, "Perform looper interpolation in an extra pass") AddOptionRTC(dropSecondaryLegsInOutput, int8_t, 1, "", 0, "Do not store secondary legs of looping track in TrackTPC") AddOptionRTC(enablePID, int8_t, 1, "", 0, "Enable PID response") AddOptionRTC(PID_useNsigma, int8_t, 1, "", 0, "Use nSigma instead of absolute distance in PID response") diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 118f0bf73a845..5d3ac212c5b54 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -224,9 +224,7 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) if (param().rec.tpc.retryRefit == 1) { runKernel(GetGridAuto(0), -1); } - if (param().rec.tpc.looperInterpolationInExtraPass == -1 ? mRec->getGPUParameters(doGPU).par_MERGER_SPLIT_LOOP_INTERPOLATION : param().rec.tpc.looperInterpolationInExtraPass) { - runKernel(GetGridAuto(0)); - } + runKernel(GetGridAuto(0)); DoDebugAndDump(RecoStep::TPCMerging, GPUChainTrackingDebugFlags::TPCMergingRefit, Merger, &GPUTPCGMMerger::DumpRefit, *mDebugFile); runKernel(GetGridAuto(0, deviceType)); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx index 87486292eb034..e4c3073f9d465 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx @@ -189,32 +189,12 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ CADEBUG(printf("\tLeg %3d Sector %2d %4sTrack Alpha %8.3f %s, X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) %28s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", (int32_t)cluster.leg, (int32_t)cluster.sector, "", prop.GetAlpha(), (CAMath::Abs(prop.GetAlpha() - clAlpha) < 0.01 ? " " : " R!"), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); // clang-format on if (allowModification && changeDirection && !noFollowCircle && !noFollowCircle2) { - bool tryFollow = lastRow != 255; - if (tryFollow) { - const GPUTPCGMTrackParam backup = *this; - const float backupAlpha = prop.GetAlpha(); - if (FollowCircle<0>(merger, prop, lastSector, lastRow, iTrk, clAlpha, xx, yy, cluster.sector, cluster.row, inFlyDirection)) { - CADEBUG(printf("Error during follow circle, resetting track!\n")); - *this = backup; - prop.SetTrack(this, backupAlpha); + if (lastRow != 255) { + if (!(merger->Param().rec.tpc.disableRefitAttachment & 4)) { + StoreAttachMirror(merger, lastSector, lastRow, iTrk, clAlpha, yy, xx, cluster.sector, cluster.row, inFlyDirection, prop.GetAlpha()); noFollowCircle = true; - tryFollow = false; } } - if (tryFollow) { - MirrorTo(prop, yy, zz, inFlyDirection, param, cluster.row, clusterState, false, cluster.sector); - lastUpdateX = mX; - lastLeg = cluster.leg; - lastSector = cluster.sector; - lastRow = 255; - N++; - resetT0 = initResetT0(); - // clang-format off - CADEBUG(printf("\n")); - CADEBUG(printf("\t%21sMirror Alpha %8.3f , X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) %28s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", "", prop.GetAlpha(), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); - // clang-format on - continue; - } } else if (allowModification && lastRow != 255 && CAMath::Abs(cluster.row - lastRow) > 1) { if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { bool dodEdx = param.dodEdxEnabled && param.rec.tpc.adddEdxSubThresholdClusters && iWay == nWays - 1 && CAMath::Abs(cluster.row - lastRow) == 2 && cluster.leg == clusters[maxN - 1].leg; @@ -269,8 +249,8 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ CADEBUG(printf(" -- MirroredY: %f --> %f", mP[0], mirrordY)); if (CAMath::Abs(yy - mP[0]) > CAMath::Abs(yy - mirrordY)) { CADEBUG(printf(" - Mirroring!!!")); - if (allowModification) { - AttachClustersMirror<0>(merger, cluster.sector, cluster.row, iTrk, yy, prop); // TODO: Never true, will always call FollowCircle above, really??? + if (allowModification && !(merger->Param().rec.tpc.disableRefitAttachment & 8)) { + StoreAttachMirror(merger, cluster.sector, cluster.row, iTrk, 0, yy, 0, -1, 0, 0, prop.GetAlpha()); } MirrorTo(prop, yy, zz, inFlyDirection, param, cluster.row, clusterState, true, cluster.sector); noFollowCircle = false; @@ -751,24 +731,15 @@ GPUdii() void GPUTPCGMTrackParam::RefitLoop(const GPUTPCGMMerger* GPUrestrict() GPUTPCGMLoopData& data = Merger->LoopData()[loopIdx]; prop.SetTrack(&data.param, data.alpha); if (data.toSector == -1) { - data.param.AttachClustersMirror<1>(Merger, data.sector, data.row, data.track, data.toY, prop, true); + data.param.AttachClustersMirror(Merger, data.sector, data.row, data.track, data.toY, prop); } else { - data.param.FollowCircle<1>(Merger, prop, data.sector, data.row, data.track, data.toAlpha, data.toX, data.toY, data.toSector, data.toRow, data.inFlyDirection, true); + data.param.FollowCircle(Merger, prop, data.sector, data.row, data.track, data.toAlpha, data.toX, data.toY, data.toSector, data.toRow, data.inFlyDirection); } } -template -GPUdic(0, 1) int32_t GPUTPCGMTrackParam::FollowCircle(const GPUTPCGMMerger* GPUrestrict() Merger, GPUTPCGMPropagator& GPUrestrict() prop, int32_t sector, int32_t iRow, int32_t iTrack, float toAlpha, float toX, float toY, int32_t toSector, int32_t toRow, bool inFlyDirection, bool phase2) +GPUdi() int32_t GPUTPCGMTrackParam::FollowCircle(const GPUTPCGMMerger* GPUrestrict() Merger, GPUTPCGMPropagator& GPUrestrict() prop, int32_t sector, int32_t iRow, int32_t iTrack, float toAlpha, float toX, float toY, int32_t toSector, int32_t toRow, bool inFlyDirection) { static constexpr float kSectAngle = 2 * M_PI / 18.f; - if (Merger->Param().rec.tpc.disableRefitAttachment & 4) { - return 1; - } - const bool inExtraPass = Merger->Param().rec.tpc.looperInterpolationInExtraPass == -1 ? GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION : Merger->Param().rec.tpc.looperInterpolationInExtraPass; - if (inExtraPass && phase2 == false) { - StoreAttachMirror(Merger, sector, iRow, iTrack, toAlpha, toY, toX, toSector, toRow, inFlyDirection, prop.GetAlpha()); - return 1; - } const GPUParam& GPUrestrict() param = Merger->Param(); bool right; float dAlpha = toAlpha - prop.GetAlpha(); @@ -862,19 +833,9 @@ GPUdic(0, 1) int32_t GPUTPCGMTrackParam::FollowCircle(const GPUTPCGMMerger* GPUr return (0); } -template -GPUdni() void GPUTPCGMTrackParam::AttachClustersMirror(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, float toY, GPUTPCGMPropagator& GPUrestrict() prop, bool phase2) +GPUdi() void GPUTPCGMTrackParam::AttachClustersMirror(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, float toY, GPUTPCGMPropagator& GPUrestrict() prop) { static constexpr float kSectAngle = 2 * M_PI / 18.f; - - if (Merger->Param().rec.tpc.disableRefitAttachment & 8) { - return; - } - const bool inExtraPass = Merger->Param().rec.tpc.looperInterpolationInExtraPass == -1 ? GPUCA_PAR_MERGER_SPLIT_LOOP_INTERPOLATION : Merger->Param().rec.tpc.looperInterpolationInExtraPass; - if (inExtraPass && phase2 == false) { - StoreAttachMirror(Merger, sector, iRow, iTrack, 0, toY, 0, -1, 0, 0, prop.GetAlpha()); - return; - } // Note that the coordinate system is rotated by 90 degree swapping X and Y! float X = mP[2] > 0 ? mP[0] : -mP[0]; float toX = mP[2] > 0 ? toY : -toY; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h index e3a5b2f7c1d01..90ff3154a3fe9 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h @@ -150,10 +150,8 @@ class GPUTPCGMTrackParam GPUd() float AttachClusters(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, GPUTPCGMPropagator& prop); // Returns uncorrectedY for later use GPUd() float AttachClusters(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, bool goodLeg, float Y, float Z); // We force to compile these twice, for RefitLoop and for Fit, for better optimization - template - GPUd() void AttachClustersMirror(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, float toY, GPUTPCGMPropagator& prop, bool phase2 = false); - template - GPUd() int32_t FollowCircle(const GPUTPCGMMerger* GPUrestrict() Merger, GPUTPCGMPropagator& prop, int32_t sector, int32_t iRow, int32_t iTrack, float toAlpha, float toX, float toY, int32_t toSector, int32_t toRow, bool inFlyDirection, bool phase2 = false); + GPUd() void AttachClustersMirror(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, float toY, GPUTPCGMPropagator& prop); + GPUd() int32_t FollowCircle(const GPUTPCGMMerger* GPUrestrict() Merger, GPUTPCGMPropagator& prop, int32_t sector, int32_t iRow, int32_t iTrack, float toAlpha, float toX, float toY, int32_t toSector, int32_t toRow, bool inFlyDirection); GPUd() void StoreAttachMirror(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t sector, int32_t iRow, int32_t iTrack, float toAlpha, float toY, float toX, int32_t toSector, int32_t toRow, bool inFlyDirection, float alpha); GPUd() void StoreOuter(gputpcgmmergertypes::GPUTPCOuterParam* outerParam, const GPUTPCGMPropagator& prop, int32_t phase); GPUd() static void RefitLoop(const GPUTPCGMMerger* GPUrestrict() Merger, int32_t loopIdx); diff --git a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx index fed4610b2f13a..5240b5ca47967 100644 --- a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx @@ -460,7 +460,6 @@ int32_t SetupReconstruction() procSet.tpcInputWithClusterRejection = 1; } recSet.tpc.disableRefitAttachment = 0xFF; - recSet.tpc.looperInterpolationInExtraPass = 0; recSet.maxTrackQPtB5 = CAMath::Min(recSet.maxTrackQPtB5, recSet.tpc.rejectQPtB5); recSet.useMatLUT = true; recAsync->SetSettings(&grp, &recSet, &procSet, &steps); diff --git a/GPU/GPUTracking/kernels.cmake b/GPU/GPUTracking/kernels.cmake index 7ebe631d86e92..e1fef5795828b 100644 --- a/GPU/GPUTracking/kernels.cmake +++ b/GPU/GPUTracking/kernels.cmake @@ -142,7 +142,6 @@ o2_gpu_kernel_add_parameter(NEIGHBOURS_FINDER_MAX_NNEIGHUP TRACKLET_SELECTOR_HITS_REG_SIZE ALTERNATE_BORDER_SORT SORT_BEFORE_FIT - MERGER_SPLIT_LOOP_INTERPOLATION NO_ATOMIC_PRECHECK COMP_GATHER_KERNEL COMP_GATHER_MODE From 636e88c447bfce9b3e1fed26a02eb929830657c7 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 23 Jul 2025 13:09:28 +0200 Subject: [PATCH 09/15] GPU TPC: Keep merged track legs as individual track segments during refit --- .../DataTypes/GPUMemorySizeScalers.h | 2 +- .../Definitions/GPUDefConstantsAndSettings.h | 2 +- GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h | 22 +- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 558 ++++++++---------- GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx | 4 +- .../display/render/GPUDisplayDraw.cxx | 14 +- GPU/GPUTracking/qa/GPUQA.cxx | 17 +- GPU/GPUTracking/qa/GPUQA.h | 3 +- 8 files changed, 284 insertions(+), 338 deletions(-) diff --git a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h index 164ecb32c26c7..ff8abdc1a491e 100644 --- a/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h +++ b/GPU/GPUTracking/DataTypes/GPUMemorySizeScalers.h @@ -47,7 +47,7 @@ struct GPUMemorySizeScalers { double tpcSectorTracksPerHit = 0.02; double tpcSectorTrackHitsPerHit = 0.8; double tpcSectorTrackHitsPerHitWithRejection = 1.0; - double tpcMergedTrackPerSectorTrack = 0.9; + double tpcMergedTrackPerSectorTrack = 1.0; double tpcMergedTrackHitPerSectorHit = 1.1; size_t tpcCompressedUnattachedHitsBase1024[3] = {900, 900, 500}; // No ratio, but integer fraction of 1024 for exact computation diff --git a/GPU/GPUTracking/Definitions/GPUDefConstantsAndSettings.h b/GPU/GPUTracking/Definitions/GPUDefConstantsAndSettings.h index 46988208256fc..04d6576f57d10 100644 --- a/GPU/GPUTracking/Definitions/GPUDefConstantsAndSettings.h +++ b/GPU/GPUTracking/Definitions/GPUDefConstantsAndSettings.h @@ -32,7 +32,7 @@ #define GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(QPTB5) (CAMath::Abs(QPTB5) > 10 ? 10 : (CAMath::Abs(QPTB5) > 5 ? 15 : 29)) // Minimum hits should depend on Pt, low Pt tracks can have few hits. 29 Hits default, 15 for < 200 mev, 10 for < 100 mev -#define GPUCA_MERGER_MAX_TRACK_CLUSTERS 1000 // Maximum number of clusters a track may have after merging +#define GPUCA_MERGER_MAX_TRACK_CLUSTERS 512 // Maximum number of clusters a track may have after merging #define GPUCA_MAXN 40 // Maximum number of neighbor hits to consider in one row in neightbors finder #define GPUCA_MIN_TRACK_PTB5_DEFAULT 0.010f // Default setting for minimum track Pt at some places (at B=0.5T) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h index 73b14ba1b2fdf..483cbc15998bc 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h @@ -45,13 +45,18 @@ class GPUTPCGMMergedTrack GPUd() bool Looper() const { return mFlags & 0x02; } GPUd() bool CSide() const { return mFlags & 0x04; } GPUd() bool CCE() const { return mFlags & 0x08; } - GPUd() bool MergedLooper() const { return mFlags & 0x10; } + GPUd() bool MergedLooperUnconnected() const { return mFlags & 0x10; } + GPUd() bool MergedLooperConnected() const { return mFlags & 0x20; } + GPUd() bool MergedLooper() const { return mFlags & 0x30; } + GPUd() int32_t PrevSegment() const { return mPrevSegment; } + GPUd() uint8_t Flags() const { return mFlags; } GPUd() void SetNClusters(int32_t v) { mNClusters = v; } GPUd() void SetNClustersFitted(int32_t v) { mNClustersFitted = v; } GPUd() void SetFirstClusterRef(int32_t v) { mFirstClusterRef = v; } GPUd() void SetParam(const GPUTPCGMTrackParam& v) { mParam = v; } GPUd() void SetAlpha(float v) { mAlpha = v; } + GPUd() void SetPrevSegment(int32_t v) { mPrevSegment = v; } GPUd() void SetOK(bool v) { if (v) { @@ -84,7 +89,7 @@ class GPUTPCGMMergedTrack mFlags &= 0xF7; } } - GPUd() void SetMergedLooper(bool v) + GPUd() void SetMergedLooperUnconnected(bool v) { if (v) { mFlags |= 0x10; @@ -92,10 +97,15 @@ class GPUTPCGMMergedTrack mFlags &= 0xEF; } } + GPUd() void SetMergedLooperConnected(bool v) + { + if (v) { + mFlags |= 0x20; + } else { + mFlags &= 0xDF; + } + } GPUd() void SetFlags(uint8_t v) { mFlags = v; } - GPUd() void SetLegs(uint8_t v) { mLegs = v; } - GPUd() uint8_t Legs() const { return mLegs; } - GPUd() uint8_t Flags() const { return mFlags; } GPUd() const gputpcgmmergertypes::GPUTPCOuterParam& OuterParam() const { return mOuterParam; } GPUd() gputpcgmmergertypes::GPUTPCOuterParam& OuterParam() { return mOuterParam; } @@ -106,11 +116,11 @@ class GPUTPCGMMergedTrack float mAlpha; //* alpha angle uint32_t mFirstClusterRef; //* index of the first track cluster in corresponding cluster arrays + int32_t mPrevSegment; //* next segment in case of looping track // TODO: Change to 8 bit uint32_t mNClusters; //* number of track clusters uint32_t mNClustersFitted; //* number of clusters used in fit uint8_t mFlags; - uint8_t mLegs; #if !defined(GPUCA_STANDALONE) ClassDefNV(GPUTPCGMMergedTrack, 0); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 338ecae4f9b95..3e2eae2e2ad6b 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -59,9 +59,6 @@ #include "SimulationDataFormat/MCCompLabel.h" #endif -static constexpr int32_t kMaxParts = 400; -static constexpr int32_t kMaxClusters = GPUCA_MERGER_MAX_TRACK_CLUSTERS; - using namespace o2::gpu; using namespace o2::tpc; using namespace gputpcgmmergertypes; @@ -98,9 +95,6 @@ struct GPUTPCGMMergerSortTracks_comp { if (a.CCE() != b.CCE()) { return a.CCE() > b.CCE(); } - if (a.Legs() != b.Legs()) { - return a.Legs() > b.Legs(); - } GPUCA_DETERMINISTIC_CODE( // clang-format off if (a.NClusters() != b.NClusters()) { return a.NClusters() > b.NClusters(); @@ -1348,8 +1342,7 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i continue; } bool celooper = (trk[0]->GetParam().GetQPt() * Param().qptB5Scaler > 1 && trk[0]->GetParam().GetQPt() * trk[1]->GetParam().GetQPt() < 0); - bool looper = trk[0]->Looper() || trk[1]->Looper() || celooper; - if (!looper && trk[0]->GetParam().GetPar(3) * trk[1]->GetParam().GetPar(3) < 0) { + if (!celooper && trk[0]->GetParam().GetPar(3) * trk[1]->GetParam().GetPar(3) < 0) { continue; } @@ -1365,7 +1358,7 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i } bool needswap = false; - if (looper) { + if (celooper) { float z0max, z1max; if (Param().par.earlyTpcTransform) { z0max = CAMath::Max(CAMath::Abs(mClustersXYZ[trk[0]->FirstClusterRef()].z), CAMath::Abs(mClustersXYZ[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].z)); @@ -1386,15 +1379,13 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i GPUCommonAlgorithm::swap(trk[0], trk[1]); } - bool reverse[2] = {false, false}; - if (looper) { - if (Param().par.earlyTpcTransform) { - reverse[0] = (mClustersXYZ[trk[0]->FirstClusterRef()].z > mClustersXYZ[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].z) ^ (trk[0]->CSide() > 0); - reverse[1] = (mClustersXYZ[trk[1]->FirstClusterRef()].z < mClustersXYZ[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].z) ^ (trk[1]->CSide() > 0); - } else { - reverse[0] = cls[mClusters[trk[0]->FirstClusterRef()].num].getTime() < cls[mClusters[trk[0]->FirstClusterRef() + trk[0]->NClusters() - 1].num].getTime(); - reverse[1] = cls[mClusters[trk[1]->FirstClusterRef()].num].getTime() > cls[mClusters[trk[1]->FirstClusterRef() + trk[1]->NClusters() - 1].num].getTime(); - } + if (celooper) { + trk[0]->SetMergedLooperConnected(true); + trk[0]->SetCCE(true); + trk[0]->SetLooper(true); + trk[1]->SetCCE(true); + trk[1]->SetLooper(true); + continue; } if (Param().par.continuousTracking) { @@ -1415,31 +1406,14 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i trk[1]->Param().TZOffset() = offset; } } - int32_t pos = newRef; - int32_t leg = -1; - int32_t lastLeg = -1; #pragma unroll for (int32_t k = 1; k >= 0; k--) { - int32_t loopstart = reverse[k] ? (trk[k]->NClusters() - 1) : 0; - int32_t loopend = reverse[k] ? -1 : (int32_t)trk[k]->NClusters(); - int32_t loopinc = reverse[k] ? -1 : 1; - for (int32_t j = loopstart; j != loopend; j += loopinc) { + for (uint32_t j = 0; j != trk[k]->NClusters(); j++) { if (Param().par.earlyTpcTransform) { mClustersXYZ[pos] = mClustersXYZ[trk[k]->FirstClusterRef() + j]; } - mClusters[pos] = mClusters[trk[k]->FirstClusterRef() + j]; - if (looper) { - if (mClusters[trk[k]->FirstClusterRef() + j].leg != lastLeg) { - leg++; - lastLeg = mClusters[trk[k]->FirstClusterRef() + j].leg; - } - mClusters[pos].leg = leg; - } - pos++; - } - if (celooper) { - lastLeg = -1; + mClusters[pos++] = mClusters[trk[k]->FirstClusterRef() + j]; } } trk[1]->SetFirstClusterRef(newRef); @@ -1449,10 +1423,6 @@ GPUd() void GPUTPCGMMerger::MergeCE(int32_t nBlocks, int32_t nThreads, int32_t i trk[1]->SetNClusters(GPUCA_MERGER_MAX_TRACK_CLUSTERS); } trk[1]->SetCCE(true); - if (looper) { - trk[1]->SetLooper(true); - trk[1]->SetLegs(leg + 1); - } trk[0]->SetNClusters(0); trk[0]->SetOK(false); } @@ -1465,32 +1435,6 @@ namespace o2::gpu::internal { namespace // anonymous { -struct GPUTPCGMMerger_CompareClusterIdsLooper { - struct clcomparestruct { - uint8_t leg; - }; - - const uint8_t leg; - const bool outwards; - const GPUTPCGMMerger::trackCluster* const cmp1; - const clcomparestruct* const cmp2; - GPUd() GPUTPCGMMerger_CompareClusterIdsLooper(uint8_t l, bool o, const GPUTPCGMMerger::trackCluster* c1, const clcomparestruct* c2) : leg(l), outwards(o), cmp1(c1), cmp2(c2) {} - GPUd() bool operator()(const int16_t aa, const int16_t bb) - { - const clcomparestruct& a = cmp2[aa]; - const clcomparestruct& b = cmp2[bb]; - const GPUTPCGMMerger::trackCluster& a1 = cmp1[aa]; - const GPUTPCGMMerger::trackCluster& b1 = cmp1[bb]; - if (a.leg != b.leg) { - return ((leg > 0) ^ (a.leg > b.leg)); - } - if (a1.row != b1.row) { - return ((a1.row > b1.row) ^ ((a.leg - leg) & 1) ^ outwards); - } - return GPUCA_DETERMINISTIC_CODE((a1.id != b1.id) ? (a1.id > b1.id) : (aa > bb), a1.id > b1.id); - } -}; - struct GPUTPCGMMerger_CompareClusterIds { const GPUTPCGMMerger::trackCluster* const mCmp; GPUd() GPUTPCGMMerger_CompareClusterIds(const GPUTPCGMMerger::trackCluster* cmp) : mCmp(cmp) {} @@ -1509,296 +1453,269 @@ struct GPUTPCGMMerger_CompareClusterIds { GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) { - GPUTPCGMSectorTrack* trackParts[kMaxParts]; + static constexpr int32_t kMaxParts = 16; + static constexpr int32_t kMaxClusters = GPUCA_MERGER_MAX_TRACK_CLUSTERS; - for (int32_t itr = iBlock * nThreads + iThread; itr < SectorTrackInfoLocalTotal(); itr += nThreads * nBlocks) { - GPUTPCGMSectorTrack& track = mSectorTrackInfos[itr]; + GPUTPCGMSectorTrack* trackParts[kMaxParts]; - if (track.PrevSegmentNeighbour() >= 0) { - continue; - } - if (track.PrevNeighbour() >= 0) { - continue; + int32_t itr = iBlock * nThreads + iThread; + GPUTPCGMSectorTrack* trbase = nullptr; + int32_t leg = 0; + int32_t lastMergedSegment = -1; + while (true) { + if (trbase) { + int32_t jtr = trbase->NextNeighbour(); + if (jtr >= 0) { + trbase = &(mSectorTrackInfos[jtr]); + if (trbase->PrevSegmentNeighbour() >= 0) { + trbase = nullptr; + } else { + trbase->SetPrevSegmentNeighbour(1000000001); + leg++; + } + } else { + trbase = nullptr; + } } - int32_t nParts = 0; - int32_t nHits = 0; - int32_t leg = 0; - GPUTPCGMSectorTrack *trbase = &track, *tr = &track; - tr->SetPrevSegmentNeighbour(1000000000); - while (true) { - if (nParts >= kMaxParts) { + + if (trbase == nullptr) { + while (itr < SectorTrackInfoLocalTotal()) { + trbase = &mSectorTrackInfos[itr]; + if (trbase->PrevSegmentNeighbour() >= 0 || trbase->PrevNeighbour() >= 0) { + itr += nThreads * nBlocks; + continue; + } break; } - if (nHits + tr->NClusters() > kMaxClusters) { + if (itr >= SectorTrackInfoLocalTotal()) { break; } - nHits += tr->NClusters(); + itr += nThreads * nBlocks; + trbase->SetPrevSegmentNeighbour(1000000000); + leg = 0; + lastMergedSegment = -1; + } - tr->SetLeg(leg); - trackParts[nParts++] = tr; - for (int32_t i = 0; i < 2; i++) { - if (tr->ExtrapolatedTrackId(i) != -1) { - if (nParts >= kMaxParts) { - break; - } - if (nHits + mSectorTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters() > kMaxClusters) { - break; - } - trackParts[nParts] = &mSectorTrackInfos[tr->ExtrapolatedTrackId(i)]; - trackParts[nParts++]->SetLeg(leg); - nHits += mSectorTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters(); + do { + int32_t nParts = 0; + int32_t nHits = 0; + + GPUTPCGMSectorTrack* tr = trbase; + while (true) { + if (nParts >= kMaxParts) { + break; } - } - int32_t jtr = tr->NextSegmentNeighbour(); - if (jtr >= 0) { - tr = &(mSectorTrackInfos[jtr]); - tr->SetPrevSegmentNeighbour(1000000002); - continue; - } - jtr = trbase->NextNeighbour(); - if (jtr >= 0) { - trbase = &(mSectorTrackInfos[jtr]); - tr = trbase; - if (tr->PrevSegmentNeighbour() >= 0) { + if (nHits + tr->NClusters() > kMaxClusters) { break; } - tr->SetPrevSegmentNeighbour(1000000001); - leg++; - continue; + nHits += tr->NClusters(); + + tr->SetLeg(leg); + trackParts[nParts++] = tr; + for (int32_t i = 0; i < 2; i++) { + if (tr->ExtrapolatedTrackId(i) != -1) { + if (nParts >= kMaxParts) { + break; + } + if (nHits + mSectorTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters() > kMaxClusters) { + break; + } + trackParts[nParts] = &mSectorTrackInfos[tr->ExtrapolatedTrackId(i)]; + trackParts[nParts++]->SetLeg(leg); + nHits += mSectorTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters(); + } + } + int32_t jtr = tr->NextSegmentNeighbour(); + if (jtr >= 0) { + tr = &(mSectorTrackInfos[jtr]); + tr->SetPrevSegmentNeighbour(1000000002); + continue; + } + break; } - break; - } - // unpack and sort clusters - if (nParts > 1 && leg == 0) { - GPUCommonAlgorithm::sort(trackParts, trackParts + nParts, [](const GPUTPCGMSectorTrack* a, const GPUTPCGMSectorTrack* b) { - GPUCA_DETERMINISTIC_CODE( // clang-format off - if (a->X() != b->X()) { + // unpack and sort clusters + if (nParts > 1 && leg == 0) { + GPUCommonAlgorithm::sort(trackParts, trackParts + nParts, [](const GPUTPCGMSectorTrack* a, const GPUTPCGMSectorTrack* b) { + GPUCA_DETERMINISTIC_CODE( // clang-format off + if (a->X() != b->X()) { + return (a->X() > b->X()); + } + if (a->Y() != b->Y()) { + return (a->Y() > b->Y()); + } + if (a->Z() != b->Z()) { + return (a->Z() > b->Z()); + } + return a->QPt() > b->QPt(); + , // !GPUCA_DETERMINISTIC_CODE return (a->X() > b->X()); - } - if (a->Y() != b->Y()) { - return (a->Y() > b->Y()); - } - if (a->Z() != b->Z()) { - return (a->Z() > b->Z()); - } - return a->QPt() > b->QPt(); - , // !GPUCA_DETERMINISTIC_CODE - return (a->X() > b->X()); - ) // clang-format on - }); - } - - if (Param().rec.tpc.dropLoopers && leg > 0) { - nParts = 1; - leg = 0; - } - - trackCluster trackClusters[kMaxClusters]; - nHits = 0; - for (int32_t ipart = 0; ipart < nParts; ipart++) { - const GPUTPCGMSectorTrack* t = trackParts[ipart]; - CADEBUG(printf("Collect Track %d Part %d QPt %f DzDs %f\n", mMemory->nMergedTracks, ipart, t->QPt(), t->DzDs())); - int32_t nTrackHits = t->NClusters(); - trackCluster* c2 = trackClusters + nHits + nTrackHits - 1; - for (int32_t i = 0; i < nTrackHits; i++, c2--) { - const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()]; - const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i]; - uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0]; - *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector(), t->Leg()}; - } - nHits += nTrackHits; - } - if (nHits < GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(track.QPt() * Param().qptB5Scaler)) { - continue; - } + ) // clang-format on + }); + } + + if (Param().rec.tpc.dropLoopers && leg > 0) { + nParts = 1; + leg = 0; + } + + trackCluster trackClusters[kMaxClusters]; + nHits = 0; + for (int32_t ipart = 0; ipart < nParts; ipart++) { + const GPUTPCGMSectorTrack* t = trackParts[ipart]; + CADEBUG(printf("Collect Track %d Part %d QPt %f DzDs %f\n", mMemory->nMergedTracks, ipart, t->QPt(), t->DzDs())); + int32_t nTrackHits = t->NClusters(); + trackCluster* c2 = trackClusters + nHits + nTrackHits - 1; + for (int32_t i = 0; i < nTrackHits; i++, c2--) { + const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()]; + const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i]; + uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0]; + *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector(), t->Leg()}; + } + nHits += nTrackHits; + } + if (nHits < GPUCA_TRACKLET_SELECTOR_MIN_HITS_B5(trbase->QPt() * Param().qptB5Scaler)) { + break; + } - int32_t ordered = leg == 0; - if (ordered) { + bool ordered = true; for (int32_t i = 1; i < nHits; i++) { if (trackClusters[i].row > trackClusters[i - 1].row || trackClusters[i].id == trackClusters[i - 1].id) { - ordered = 0; + ordered = false; break; } } - } - int32_t firstTrackIndex = 0; - int32_t lastTrackIndex = nParts - 1; - if (ordered == 0) { - int32_t nTmpHits = 0; - trackCluster trackClustersUnsorted[kMaxClusters]; - int16_t clusterIndices[kMaxClusters]; - for (int32_t i = 0; i < nHits; i++) { - trackClustersUnsorted[i] = trackClusters[i]; - clusterIndices[i] = i; - } + int32_t firstTrackIndex = 0; + int32_t lastTrackIndex = nParts - 1; + if (ordered == 0) { + int32_t nTmpHits = 0; + trackCluster trackClustersUnsorted[kMaxClusters]; + int16_t clusterIndices[kMaxClusters]; + for (int32_t i = 0; i < nHits; i++) { + trackClustersUnsorted[i] = trackClusters[i]; + clusterIndices[i] = i; + } + + GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIds(trackClusters)); - if (leg > 0) { - // Find QPt and DzDs for the segment closest to the vertex, if low/mid Pt - float baseZT = 1e9; - uint8_t baseLeg = 0; + nTmpHits = 0; + firstTrackIndex = lastTrackIndex = -1; for (int32_t i = 0; i < nParts; i++) { - if (trackParts[i]->Leg() == 0 || trackParts[i]->Leg() == leg) { - float zt; - if (Param().par.earlyTpcTransform) { - zt = CAMath::Min(CAMath::Abs(trackParts[i]->ClusterZT0()), CAMath::Abs(trackParts[i]->ClusterZTN())); - } else { - zt = -trackParts[i]->MinClusterZT(); // Negative time ~ smallest z, to behave the same way // TODO: Check all these min / max ZT - } - if (zt < baseZT) { - baseZT = zt; - baseLeg = trackParts[i]->Leg(); - } + nTmpHits += trackParts[i]->NClusters(); + if (nTmpHits > clusterIndices[0] && firstTrackIndex == -1) { + firstTrackIndex = i; } - } - int32_t iLongest = 1e9; - int32_t length = 0; - for (int32_t i = (baseLeg ? (nParts - 1) : 0); baseLeg ? (i >= 0) : (i < nParts); baseLeg ? i-- : i++) { - if (trackParts[i]->Leg() != baseLeg) { - break; - } - if (trackParts[i]->OrigTrack()->NHits() > length) { - iLongest = i; - length = trackParts[i]->OrigTrack()->NHits(); + if (nTmpHits > clusterIndices[nHits - 1] && lastTrackIndex == -1) { + lastTrackIndex = i; } } - bool outwards; - if (Param().par.earlyTpcTransform) { - outwards = (trackParts[iLongest]->ClusterZT0() > trackParts[iLongest]->ClusterZTN()) ^ trackParts[iLongest]->CSide(); - } else { - outwards = trackParts[iLongest]->ClusterZT0() < trackParts[iLongest]->ClusterZTN(); - } - GPUTPCGMMerger_CompareClusterIdsLooper::clcomparestruct clusterSort[kMaxClusters]; - for (int32_t iPart = 0; iPart < nParts; iPart++) { - const GPUTPCGMSectorTrack* t = trackParts[iPart]; - int32_t nTrackHits = t->NClusters(); - for (int32_t j = 0; j < nTrackHits; j++) { - int32_t i = nTmpHits + j; - clusterSort[i].leg = t->Leg(); + + int32_t nFilteredHits = 0; + int32_t indPrev = -1; + for (int32_t i = 0; i < nHits; i++) { + int32_t ind = clusterIndices[i]; + if (indPrev >= 0 && trackClustersUnsorted[ind].id == trackClustersUnsorted[indPrev].id) { + continue; } - nTmpHits += nTrackHits; + indPrev = ind; + trackClusters[nFilteredHits] = trackClustersUnsorted[ind]; + nFilteredHits++; } - - GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIdsLooper(baseLeg, outwards, trackClusters, clusterSort)); - } else { - GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIds(trackClusters)); + nHits = nFilteredHits; } - nTmpHits = 0; - firstTrackIndex = lastTrackIndex = -1; - for (int32_t i = 0; i < nParts; i++) { - nTmpHits += trackParts[i]->NClusters(); - if (nTmpHits > clusterIndices[0] && firstTrackIndex == -1) { - firstTrackIndex = i; - } - if (nTmpHits > clusterIndices[nHits - 1] && lastTrackIndex == -1) { - lastTrackIndex = i; - } + + const uint32_t iMergedTrackFirstCluster = CAMath::AtomicAdd(&mMemory->nMergedTrackClusters, (uint32_t)nHits); + if (iMergedTrackFirstCluster + nHits > mNMaxMergedTrackClusters) { + raiseError(GPUErrors::ERROR_MERGER_HIT_OVERFLOW, iMergedTrackFirstCluster, mNMaxMergedTrackClusters); + CAMath::AtomicExch(&mMemory->nMergedTrackClusters, mNMaxMergedTrackClusters); + break; } - int32_t nFilteredHits = 0; - int32_t indPrev = -1; + GPUTPCGMMergedTrackHit* const cl = mClusters + iMergedTrackFirstCluster; + for (int32_t i = 0; i < nHits; i++) { - int32_t ind = clusterIndices[i]; - if (indPrev >= 0 && trackClustersUnsorted[ind].id == trackClustersUnsorted[indPrev].id) { - continue; + uint8_t state; + if (Param().par.earlyTpcTransform) { + const GPUTPCClusterData& c = GetConstantMem()->tpcTrackers[trackClusters[i].sector].ClusterData()[trackClusters[i].id - GetConstantMem()->tpcTrackers[trackClusters[i].sector].Data().ClusterIdOffset()]; + GPUTPCGMMergedTrackHitXYZ* const clXYZ = mClustersXYZ + iMergedTrackFirstCluster; + clXYZ[i].x = c.x; + clXYZ[i].y = c.y; + clXYZ[i].z = c.z; + clXYZ[i].amp = c.amp; + state = c.flags; + } else { + const ClusterNative& c = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[trackClusters[i].id]; + state = c.getFlags(); } - indPrev = ind; - trackClusters[nFilteredHits] = trackClustersUnsorted[ind]; - nFilteredHits++; + cl[i].state = state & GPUTPCGMMergedTrackHit::clustererAndSharedFlags; // Only allow edge, deconvoluted, and shared flags + cl[i].row = trackClusters[i].row; + cl[i].num = trackClusters[i].id; + cl[i].sector = trackClusters[i].sector; + cl[i].leg = trackClusters[i].leg; } - nHits = nFilteredHits; - } - - const uint32_t iMergedTrackFirstCluster = CAMath::AtomicAdd(&mMemory->nMergedTrackClusters, (uint32_t)nHits); - if (iMergedTrackFirstCluster + nHits > mNMaxMergedTrackClusters) { - raiseError(GPUErrors::ERROR_MERGER_HIT_OVERFLOW, iMergedTrackFirstCluster, mNMaxMergedTrackClusters); - CAMath::AtomicExch(&mMemory->nMergedTrackClusters, mNMaxMergedTrackClusters); - continue; - } - GPUTPCGMMergedTrackHit* const cl = mClusters + iMergedTrackFirstCluster; + uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nMergedTracks, 1u); + if (iOutputTrack >= mNMaxTracks) { + raiseError(GPUErrors::ERROR_MERGER_TRACK_OVERFLOW, iOutputTrack, mNMaxTracks); + CAMath::AtomicExch(&mMemory->nMergedTracks, mNMaxTracks); + break; + } - for (int32_t i = 0; i < nHits; i++) { - uint8_t state; - if (Param().par.earlyTpcTransform) { - const GPUTPCClusterData& c = GetConstantMem()->tpcTrackers[trackClusters[i].sector].ClusterData()[trackClusters[i].id - GetConstantMem()->tpcTrackers[trackClusters[i].sector].Data().ClusterIdOffset()]; - GPUTPCGMMergedTrackHitXYZ* const clXYZ = mClustersXYZ + iMergedTrackFirstCluster; - clXYZ[i].x = c.x; - clXYZ[i].y = c.y; - clXYZ[i].z = c.z; - clXYZ[i].amp = c.amp; - state = c.flags; + GPUTPCGMMergedTrack& mergedTrack = mMergedTracks[iOutputTrack]; + mergedTrack.SetFlags(0); + mergedTrack.SetOK(true); + mergedTrack.SetLooper(leg > 0 || trbase->NextNeighbour() >= 0); + mergedTrack.SetNClusters(nHits); + mergedTrack.SetFirstClusterRef(iMergedTrackFirstCluster); + GPUTPCGMTrackParam& p1 = mergedTrack.Param(); + const GPUTPCGMSectorTrack& p2 = *trackParts[firstTrackIndex]; + mergedTrack.SetCSide(p2.CSide()); + mergedTrack.SetMergedLooperConnected(leg > 0); + mergedTrack.SetPrevSegment(lastMergedSegment); + lastMergedSegment = iOutputTrack; + + GPUTPCGMBorderTrack b; + const float toX = Param().par.earlyTpcTransform ? mClustersXYZ[iMergedTrackFirstCluster].x : GPUTPCGeometry::Row2X(cl[0].row); + if (p2.TransportToX(this, toX, Param().bzCLight, b, GPUCA_MAX_SIN_PHI, false)) { + p1.X() = toX; + p1.Y() = b.Par()[0]; + p1.Z() = b.Par()[1]; + p1.SinPhi() = b.Par()[2]; } else { - const ClusterNative& c = GetConstantMem()->ioPtrs.clustersNative->clustersLinear[trackClusters[i].id]; - state = c.getFlags(); - } - cl[i].state = state & GPUTPCGMMergedTrackHit::clustererAndSharedFlags; // Only allow edge, deconvoluted, and shared flags - cl[i].row = trackClusters[i].row; - cl[i].num = trackClusters[i].id; - cl[i].sector = trackClusters[i].sector; - cl[i].leg = trackClusters[i].leg; - } - - uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nMergedTracks, 1u); - if (iOutputTrack >= mNMaxTracks) { - raiseError(GPUErrors::ERROR_MERGER_TRACK_OVERFLOW, iOutputTrack, mNMaxTracks); - CAMath::AtomicExch(&mMemory->nMergedTracks, mNMaxTracks); - continue; - } - - GPUTPCGMMergedTrack& mergedTrack = mMergedTracks[iOutputTrack]; - - mergedTrack.SetFlags(0); - mergedTrack.SetOK(1); - mergedTrack.SetLooper(leg > 0); - mergedTrack.SetLegs(leg); - mergedTrack.SetNClusters(nHits); - mergedTrack.SetFirstClusterRef(iMergedTrackFirstCluster); - GPUTPCGMTrackParam& p1 = mergedTrack.Param(); - const GPUTPCGMSectorTrack& p2 = *trackParts[firstTrackIndex]; - mergedTrack.SetCSide(p2.CSide()); - - GPUTPCGMBorderTrack b; - const float toX = Param().par.earlyTpcTransform ? mClustersXYZ[iMergedTrackFirstCluster].x : GPUTPCGeometry::Row2X(cl[0].row); - if (p2.TransportToX(this, toX, Param().bzCLight, b, GPUCA_MAX_SIN_PHI, false)) { - p1.X() = toX; - p1.Y() = b.Par()[0]; - p1.Z() = b.Par()[1]; - p1.SinPhi() = b.Par()[2]; - } else { - p1.X() = p2.X(); - p1.Y() = p2.Y(); - p1.Z() = p2.Z(); - p1.SinPhi() = p2.SinPhi(); - } - p1.TZOffset() = p2.TZOffset(); - p1.DzDs() = p2.DzDs(); - p1.QPt() = p2.QPt(); - mergedTrack.SetAlpha(p2.Alpha()); - if (CAMath::Abs(Param().polynomialField.GetNominalBz()) < (gpu_common_constants::kZeroFieldCut * gpu_common_constants::kCLight)) { - p1.QPt() = 100.f / Param().rec.bz0Pt10MeV; - } + p1.X() = p2.X(); + p1.Y() = p2.Y(); + p1.Z() = p2.Z(); + p1.SinPhi() = p2.SinPhi(); + } + p1.TZOffset() = p2.TZOffset(); + p1.DzDs() = p2.DzDs(); + p1.QPt() = p2.QPt(); + mergedTrack.SetAlpha(p2.Alpha()); + if (CAMath::Abs(Param().polynomialField.GetNominalBz()) < (gpu_common_constants::kZeroFieldCut * gpu_common_constants::kCLight)) { + p1.QPt() = 100.f / Param().rec.bz0Pt10MeV; + } - // if (nParts > 1) printf("Merged %d: QPt %f %d parts %d hits\n", mMemory->nMergedTracks, p1.QPt(), nParts, nHits); + // if (nParts > 1) printf("Merged %d: QPt %f %d parts %d hits\n", mMemory->nMergedTracks, p1.QPt(), nParts, nHits); - /*if (GPUQA::QAAvailable() && mRec->GetQA() && mRec->GetQA()->SuppressTrack(mMemory->nMergedTracks)) - { - mergedTrack.SetOK(0); - mergedTrack.SetNClusters(0); - } - if (mergedTrack.NClusters() && mergedTrack.OK()) */ - if (Param().rec.tpc.mergeCE) { - bool CEside; - if (Param().par.earlyTpcTransform) { - const GPUTPCGMMergedTrackHitXYZ* const clXYZ = mClustersXYZ + iMergedTrackFirstCluster; - CEside = (mergedTrack.CSide() != 0) ^ (clXYZ[0].z > clXYZ[nHits - 1].z); - } else { - auto& cls = mConstantMem->ioPtrs.clustersNative->clustersLinear; - CEside = cls[cl[0].num].getTime() < cls[cl[nHits - 1].num].getTime(); + /*if (GPUQA::QAAvailable() && mRec->GetQA() && mRec->GetQA()->SuppressTrack(mMemory->nMergedTracks)) + { + mergedTrack.SetOK(0); + mergedTrack.SetNClusters(0); } - MergeCEFill(trackParts[CEside ? lastTrackIndex : firstTrackIndex], cl[CEside ? (nHits - 1) : 0], Param().par.earlyTpcTransform ? &(mClustersXYZ + iMergedTrackFirstCluster)[CEside ? (nHits - 1) : 0] : nullptr, iOutputTrack); - } - } // itr + if (mergedTrack.NClusters() && mergedTrack.OK()) */ + if (Param().rec.tpc.mergeCE) { + bool CEside; + if (Param().par.earlyTpcTransform) { + const GPUTPCGMMergedTrackHitXYZ* const clXYZ = mClustersXYZ + iMergedTrackFirstCluster; + CEside = (mergedTrack.CSide() != 0) ^ (clXYZ[0].z > clXYZ[nHits - 1].z); + } else { + auto& cls = mConstantMem->ioPtrs.clustersNative->clustersLinear; + CEside = cls[cl[0].num].getTime() < cls[cl[nHits - 1].num].getTime(); + } + MergeCEFill(trackParts[CEside ? lastTrackIndex : firstTrackIndex], cl[CEside ? (nHits - 1) : 0], Param().par.earlyTpcTransform ? &(mClustersXYZ + iMergedTrackFirstCluster)[CEside ? (nHits - 1) : 0] : nullptr, iOutputTrack); + } + } while (false); + } } GPUd() void GPUTPCGMMerger::SortTracksPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) @@ -1911,6 +1828,7 @@ GPUd() void GPUTPCGMMerger::Finalize2(int32_t nBlocks, int32_t nThreads, int32_t GPUd() void GPUTPCGMMerger::MergeLoopersInit(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) { + return; // FIXME: !!!! const float lowPtThresh = Param().rec.tpc.rejectQPtB5 * 1.1f; // Might need to merge tracks above the threshold with parts below the threshold for (uint32_t i = get_global_id(0); i < mMemory->nMergedTracks; i += get_global_size(0)) { const auto& trk = mMergedTracks[i]; @@ -2057,9 +1975,9 @@ GPUd() void GPUTPCGMMerger::MergeLoopersMain(int32_t nBlocks, int32_t nThreads, }*/ #endif if (EQ) { - mMergedTracks[params[j].id].SetMergedLooper(true); + mMergedTracks[params[j].id].SetMergedLooperUnconnected(true); if (CAMath::Abs(param2.GetQPt() * Param().qptB5Scaler) >= Param().rec.tpc.rejectQPtB5) { - mMergedTracks[params[i].id].SetMergedLooper(true); + mMergedTracks[params[i].id].SetMergedLooperUnconnected(true); } } } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx index 9c924e74ec519..90f2fce5cdd2e 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx @@ -205,7 +205,7 @@ void GPUTPCGMMerger::DumpRefit(std::ostream& out) const 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() << " dEdx " << (trk.OK() && Param().dodEdxEnabled ? mMergedTracksdEdx[i].dEdxTotTPC : -1.f) << "/" << (trk.OK() && Param().dodEdxEnabled ? mMergedTracksdEdx[i].dEdxMaxTPC : -1.f) << " Outer " << po.P[0] << "/" << po.P[1] << "/" << po.P[2] << "/" << po.P[3] << "/" << po.P[4] - << " NFitted " << trk.NClustersFitted() << " legs " << (int)trk.Legs() << " flags " << (int)trk.Flags() << "\n"; + << " NFitted " << trk.NClustersFitted() << " flags " << (int)trk.Flags() << "\n"; } out << std::setprecision(ss); } @@ -217,7 +217,7 @@ void GPUTPCGMMerger::DumpLoopers(std::ostream& out) const if (i && i % 100 == 0) { out << "\n"; } - out << (int)mMergedTracks[i].MergedLooper() << " "; + out << (int)mMergedTracks[i].MergedLooperUnconnected() << " "; } out << "\n"; } diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index f50e6590eb16d..2695a732b17c7 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -464,6 +464,13 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp } else { if (!drawing) { startCountInner = mVertexBuffer[iSector].size(); + if constexpr (std::is_same_v) { + if (k == 0 && track->PrevSegment() >= 0) { + const auto& prevtrk = mIOPtrs->mergedTracks[track->PrevSegment()]; + int32_t prevcid = mIOPtrs->mergedTrackHits[prevtrk.FirstClusterRef() + prevtrk.NClusters() - 1].num; + drawPointLinestrip(iSector, prevcid, tFINALTRACK, separateExtrapolatedTracksLimit); + } + } if (lastCluster != -1 && (!mCfgH.splitCETracks || lastSide == (mGlobalPos[cid].z < 0))) { int32_t lastcid; if constexpr (std::is_same_v) { @@ -512,6 +519,11 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp if (lastCluster == -1) { continue; } + if constexpr (std::is_same_v) { + if (track->MergedLooperConnected()) { + continue; + } + } } size_t startCountInner = mVertexBuffer[iSector].size(); @@ -607,7 +619,7 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp if ((inFlyDirection == 0 && x < 0) || (inFlyDirection && x * x + trkParam.Y() * trkParam.Y() > (iMC ? (450 * 450) : (300 * 300)))) { break; } - if (fabsf(trkParam.Z() + ZOffset) > mMaxClusterZ + (iMC ? 0 : 0)) { + if (fabsf(trkParam.Z() + ZOffset) > mMaxClusterZ) { break; } if (fabsf(trkParam.Z() - z0) > (iMC ? GPUTPCGeometry::TPCLength() : GPUTPCGeometry::TPCLength())) { diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx index 673f63793939d..f6ce56df6492a 100644 --- a/GPU/GPUTracking/qa/GPUQA.cxx +++ b/GPU/GPUTracking/qa/GPUQA.cxx @@ -100,13 +100,15 @@ using namespace o2::gpu; float qpt = 0; \ bool lowPt = false; \ [[maybe_unused]] bool mev200 = false; \ - bool mergedLooper = false; \ + bool mergedLooperUnconnected = false; \ + bool mergedLooperConnected = false; \ int32_t id = attach & gputpcgmmergertypes::attachTrackMask; \ if (!unattached) { \ qpt = fabsf(mTracking->mIOPtrs.mergedTracks[id].GetParam().GetQPt()); \ lowPt = qpt * mTracking->GetParam().qptB5Scaler > mTracking->GetParam().rec.tpc.rejectQPtB5; \ mev200 = qpt > 5; \ - mergedLooper = mTracking->mIOPtrs.mergedTracks[id].MergedLooper(); \ + mergedLooperUnconnected = mTracking->mIOPtrs.mergedTracks[id].MergedLooperUnconnected(); \ + mergedLooperConnected = mTracking->mIOPtrs.mergedTracks[id].MergedLooperConnected(); \ } \ bool physics = false, protect = false; \ CHECK_CLUSTER_STATE_INIT_LEG_BY_MC(); @@ -118,15 +120,17 @@ using namespace o2::gpu; } \ if (lowPt) { \ mClusterCounts.nLowPt++; \ - } else if (mergedLooper) { \ - mClusterCounts.nMergedLooper++; \ + } else if (mergedLooperUnconnected) { \ + mClusterCounts.nMergedLooperUnconnected++; \ + } else if (mergedLooperConnected) { \ + mClusterCounts.nMergedLooperConnected++; \ } else { \ GPUTPCClusterRejection::GetProtectionStatus(attach, physics, protect, &mClusterCounts, &mev200); \ } #define CHECK_CLUSTER_STATE_NOCOUNT() \ CHECK_CLUSTER_STATE_INIT() \ - if (!lowPt && !mergedLooper) { \ + if (!lowPt && !mergedLooperUnconnected && !mergedLooperConnected) { \ GPUTPCClusterRejection::GetProtectionStatus(attach, physics, protect); \ } @@ -2886,7 +2890,8 @@ int32_t GPUQA::DoClusterCounts(uint64_t* attachClusterCounts, int32_t mode) PrintClusterCount(mode, num, "Removed (Strategy B)", mClusterCounts.nTotal - mClusterCounts.nProt, mClusterCounts.nTotal); } - PrintClusterCount(mode, num, "Merged Loopers (Afterburner)", mClusterCounts.nMergedLooper, mClusterCounts.nTotal); + PrintClusterCount(mode, num, "Merged Loopers (Track Merging)", mClusterCounts.nMergedLooperConnected, mClusterCounts.nTotal); + PrintClusterCount(mode, num, "Merged Loopers (Afterburner)", mClusterCounts.nMergedLooperUnconnected, mClusterCounts.nTotal); PrintClusterCount(mode, num, "High Inclination Angle", mClusterCounts.nHighIncl, mClusterCounts.nTotal); PrintClusterCount(mode, num, "Rejected", mClusterCounts.nRejected, mClusterCounts.nTotal); PrintClusterCount(mode, num, "Tube (> 200 MeV)", mClusterCounts.nTube, mClusterCounts.nTotal); diff --git a/GPU/GPUTracking/qa/GPUQA.h b/GPU/GPUTracking/qa/GPUQA.h index e587b15f68d80..d6889ad7d3d02 100644 --- a/GPU/GPUTracking/qa/GPUQA.h +++ b/GPU/GPUTracking/qa/GPUQA.h @@ -291,7 +291,8 @@ class GPUQA TLegend* mLClust[N_CLS_TYPE]; struct counts_t { - int64_t nRejected = 0, nTube = 0, nTube200 = 0, nLoopers = 0, nLowPt = 0, n200MeV = 0, nPhysics = 0, nProt = 0, nUnattached = 0, nTotal = 0, nHighIncl = 0, nAbove400 = 0, nFakeRemove400 = 0, nFullFakeRemove400 = 0, nBelow40 = 0, nFakeProtect40 = 0, nMergedLooper = 0, nCorrectlyAttachedNormalized = 0, nCorrectlyAttachedNormalizedNonFake = 0; + int64_t nRejected = 0, nTube = 0, nTube200 = 0, nLoopers = 0, nLowPt = 0, n200MeV = 0, nPhysics = 0, nProt = 0, nUnattached = 0, nTotal = 0, nHighIncl = 0, nAbove400 = 0, nFakeRemove400 = 0, nFullFakeRemove400 = 0, nBelow40 = 0, nFakeProtect40 = 0; + int64_t nMergedLooperConnected = 0, nMergedLooperUnconnected = 0, nCorrectlyAttachedNormalized = 0, nCorrectlyAttachedNormalizedNonFake = 0; double nUnaccessible = 0; } mClusterCounts; From b8a72f7a4280494781a8606aad2f5c9ac563ad43 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 24 Jul 2025 11:47:33 +0200 Subject: [PATCH 10/15] GPU TPC: Order legs in descending way and store leg id per track not cluster --- .../GPUTPCCompressionKernels.cxx | 8 +-- .../DataTypes/GPUTPCGMMergedTrackHit.h | 2 +- GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h | 3 ++ GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 26 ++++------ GPU/GPUTracking/Merger/GPUTPCGMMerger.h | 1 - GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx | 10 ++-- GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx | 52 +++++-------------- GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h | 2 +- GPU/GPUTracking/Refit/GPUTrackingRefit.cxx | 9 ---- GPU/GPUTracking/TRDTracking/GPUTRDTracker.h | 2 +- .../display/render/GPUDisplayDraw.cxx | 22 +++++--- GPU/GPUTracking/qa/GPUQA.cxx | 2 +- 12 files changed, 52 insertions(+), 87 deletions(-) diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx index 73b195e8f4fe4..5503eeb30cdd6 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx @@ -32,7 +32,6 @@ GPUdii() void GPUTPCCompressionKernels::ThreadnStoredTracks, 1u); compressor.mAttachedClusterFirstIndex[myTrack] = trk.FirstClusterRef(); - lastLeg = hit.leg; c.qPtA[myTrack] = qpt; c.rowA[myTrack] = hit.row; c.sliceA[myTrack] = hit.sector; @@ -114,12 +109,11 @@ GPUdii() void GPUTPCCompressionKernels::Thread= 0) { continue; } - int32_t leg = 0; GPUTPCGMSectorTrack *trbase = &track, *tr = &track; while (true) { int32_t iTrk = tr - mSectorTrackInfos; @@ -200,7 +199,6 @@ void GPUTPCGMMerger::CheckMergedTracks() if (tr->PrevSegmentNeighbour() >= 0) { break; } - leg++; continue; } break; @@ -1463,7 +1461,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread int32_t leg = 0; int32_t lastMergedSegment = -1; while (true) { - if (trbase) { + if (trbase && !Param().rec.tpc.dropLoopers) { int32_t jtr = trbase->NextNeighbour(); if (jtr >= 0) { trbase = &(mSectorTrackInfos[jtr]); @@ -1471,7 +1469,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread trbase = nullptr; } else { trbase->SetPrevSegmentNeighbour(1000000001); - leg++; + leg--; } } else { trbase = nullptr; @@ -1492,7 +1490,12 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread } itr += nThreads * nBlocks; trbase->SetPrevSegmentNeighbour(1000000000); + int32_t jtr = trbase->NextNeighbour(); leg = 0; + while (jtr >= 0) { + leg++; + jtr = mSectorTrackInfos[jtr].NextNeighbour(); + } lastMergedSegment = -1; } @@ -1535,7 +1538,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread } // unpack and sort clusters - if (nParts > 1 && leg == 0) { + if (nParts > 1) { GPUCommonAlgorithm::sort(trackParts, trackParts + nParts, [](const GPUTPCGMSectorTrack* a, const GPUTPCGMSectorTrack* b) { GPUCA_DETERMINISTIC_CODE( // clang-format off if (a->X() != b->X()) { @@ -1554,11 +1557,6 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread }); } - if (Param().rec.tpc.dropLoopers && leg > 0) { - nParts = 1; - leg = 0; - } - trackCluster trackClusters[kMaxClusters]; nHits = 0; for (int32_t ipart = 0; ipart < nParts; ipart++) { @@ -1570,7 +1568,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[t->Sector()]; const GPUTPCHitId& ic = trk.TrackHits()[t->OrigTrack()->FirstHitID() + i]; uint32_t id = trk.Data().ClusterDataIndex(trk.Data().Row(ic.RowIndex()), ic.HitIndex()) + GetConstantMem()->ioPtrs.clustersNative->clusterOffset[t->Sector()][0]; - *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector(), t->Leg()}; + *c2 = trackCluster{id, (uint8_t)ic.RowIndex(), t->Sector()}; } nHits += nTrackHits; } @@ -1651,7 +1649,6 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread cl[i].row = trackClusters[i].row; cl[i].num = trackClusters[i].id; cl[i].sector = trackClusters[i].sector; - cl[i].leg = trackClusters[i].leg; } uint32_t iOutputTrack = CAMath::AtomicAdd(&mMemory->nMergedTracks, 1u); @@ -1664,7 +1661,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread GPUTPCGMMergedTrack& mergedTrack = mMergedTracks[iOutputTrack]; mergedTrack.SetFlags(0); mergedTrack.SetOK(true); - mergedTrack.SetLooper(leg > 0 || trbase->NextNeighbour() >= 0); + mergedTrack.SetLooper(leg > 0 || lastMergedSegment >= 0); mergedTrack.SetNClusters(nHits); mergedTrack.SetFirstClusterRef(iMergedTrackFirstCluster); GPUTPCGMTrackParam& p1 = mergedTrack.Param(); @@ -1799,7 +1796,6 @@ GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t if (!trk.OK() || trk.NClusters() == 0) { continue; } - uint8_t goodLeg = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].leg; for (uint32_t j = 0; j < trk.NClusters(); j++) { int32_t id = mClusters[trk.FirstClusterRef() + j].num; uint32_t weight = mTrackOrderAttach[i] | attachAttached; @@ -1809,7 +1805,7 @@ GPUd() void GPUTPCGMMerger::Finalize1(int32_t nBlocks, int32_t nThreads, int32_t } else if (clusterState & GPUTPCGMMergedTrackHit::flagHighIncl) { weight |= attachHighIncl; } - if (mClusters[trk.FirstClusterRef() + j].leg == goodLeg) { + if (trk.Leg() == 0) { weight |= attachGoodLeg; } CAMath::AtomicMax(&mClusterAttachment[id], weight); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index 54a541ebe0fd6..76f3f3cdcba08 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -82,7 +82,6 @@ class GPUTPCGMMerger : public GPUProcessor uint32_t id; uint8_t row; uint8_t sector; - uint8_t leg; }; struct tmpSort { diff --git a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx index 74a8df388d163..b10b1d0510fd7 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMO2Output.cxx @@ -65,14 +65,15 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlock if (!tracks[i].OK()) { continue; } + if (merger.Param().rec.tpc.dropSecondaryLegsInOutput && tracks[i].MergedLooper()) { + continue; + } + uint32_t nCl = 0; for (uint32_t j = 0; j < tracks[i].NClusters(); j++) { if ((trackClusters[tracks[i].FirstClusterRef() + j].state & flagsReject) || (merger.ClusterAttachment()[trackClusters[tracks[i].FirstClusterRef() + j].num] & flagsRequired) != flagsRequired) { continue; } - if (merger.Param().rec.tpc.dropSecondaryLegsInOutput && trackClusters[tracks[i].FirstClusterRef() + j].leg != trackClusters[tracks[i].FirstClusterRef() + tracks[i].NClusters() - 1].leg) { - continue; - } nCl++; } if (nCl == 0) { @@ -192,9 +193,6 @@ GPUdii() void GPUTPCGMO2Output::Thread(int32_t nBlocks if ((trackClusters[tracks[i].FirstClusterRef() + j].state & flagsReject) || (merger.ClusterAttachment()[trackClusters[tracks[i].FirstClusterRef() + j].num] & flagsRequired) != flagsRequired) { continue; } - if (merger.Param().rec.tpc.dropSecondaryLegsInOutput && trackClusters[tracks[i].FirstClusterRef() + j].leg != trackClusters[tracks[i].FirstClusterRef() + tracks[i].NClusters() - 1].leg) { - continue; - } int32_t clusterIdGlobal = trackClusters[tracks[i].FirstClusterRef() + j].num; int32_t sector = trackClusters[tracks[i].FirstClusterRef() + j].sector; int32_t globalRow = trackClusters[tracks[i].FirstClusterRef() + j].row; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx index e4c3073f9d465..db357ff034e90 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx @@ -54,7 +54,7 @@ using namespace o2::gpu; using namespace o2::tpc; -GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_t iTrk, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, GPUTPCGMMergedTrackHitXYZ* GPUrestrict() clustersXYZ, int32_t& GPUrestrict() N, int32_t& GPUrestrict() NTolerated, float& GPUrestrict() Alpha, int32_t attempt, float maxSinPhi, gputpcgmmergertypes::GPUTPCOuterParam* GPUrestrict() outerParam) +GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_t iTrk, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, GPUTPCGMMergedTrackHitXYZ* GPUrestrict() clustersXYZ, int32_t& GPUrestrict() N, int32_t& GPUrestrict() NTolerated, float& GPUrestrict() Alpha, int32_t attempt, float maxSinPhi, gputpcgmmergertypes::GPUTPCOuterParam* GPUrestrict() outerParam, int8_t leg) { static constexpr float kDeg2Rad = M_PI / 180.f; CADEBUG(static constexpr float kSectAngle = 2 * M_PI / 18.f); @@ -83,22 +83,15 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ float lastUpdateX = -1.f; uint8_t lastRow = 255; uint8_t lastSector = 255; - uint8_t storeOuter = 0; for (int32_t iWay = 0; iWay < nWays; iWay++) { int32_t nMissed = 0, nMissed2 = 0; float sumInvSqrtCharge = 0.f; int32_t nAvgCharge = 0; - if (iWay && storeOuter != 255 && param.rec.tpc.nWaysOuter && outerParam) { - storeOuter = 0; + if (iWay && param.rec.tpc.nWaysOuter && outerParam) { if (iWay == nWays - 1) { StoreOuter(outerParam, prop, 0); - if (merger->MergedTracks()[iTrk].Looper()) { - storeOuter = 1; - } - } else if (iWay == nWays - 2 && merger->MergedTracks()[iTrk].Looper()) { - storeOuter = 2; } } @@ -117,8 +110,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ N = 0; lastUpdateX = -1; - const bool inFlyDirection = iWay & 1; - uint8_t lastLeg = clusters[ihitStart].leg; + const bool inFlyDirection = (leg & 1); const int32_t wayDirection = (iWay & 1) ? -1 : 1; bool noFollowCircle = false, noFollowCircle2 = false; @@ -130,15 +122,6 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ noFollowCircle2 = true; } - if (storeOuter == 2 && clusters[ihit].leg == clusters[maxN - 1].leg - 1) { - if (lastLeg == clusters[maxN - 1].leg) { - StoreOuter(outerParam, prop, 1); - storeOuter = 255; - } else { - storeOuter = 0; - } - } - if ((param.rec.tpc.trackFitRejectMode > 0 && nMissed >= param.rec.tpc.trackFitRejectMode) || nMissed2 >= param.rec.tpc.trackFitMaxRowMissedHard || clusters[ihit].state & GPUTPCGMMergedTrackHit::flagReject) { CADEBUG(printf("\tSkipping hit, %d hits rejected, flag %X\n", nMissed, (int32_t)clusters[ihit].state)); if (iWay + 2 >= nWays && !(clusters[ihit].state & GPUTPCGMMergedTrackHit::flagReject)) { @@ -183,12 +166,10 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ const auto& cluster = clusters[ihit]; - bool changeDirection = (cluster.leg - lastLeg) & 1; // clang-format off - CADEBUG(if (changeDirection) printf("\t\tChange direction\n")); - CADEBUG(printf("\tLeg %3d Sector %2d %4sTrack Alpha %8.3f %s, X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) %28s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", (int32_t)cluster.leg, (int32_t)cluster.sector, "", prop.GetAlpha(), (CAMath::Abs(prop.GetAlpha() - clAlpha) < 0.01 ? " " : " R!"), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); + CADEBUG(printf("\tSector %2d %4sTrack Alpha %8.3f %s, X %8.3f - Y %8.3f, Z %8.3f - QPt %7.2f (%7.2f), SP %5.2f (%5.2f) %28s --- Cov sY %8.3f sZ %8.3f sSP %8.3f sPt %8.3f - YPt %8.3f\n", (int32_t)cluster.sector, "", prop.GetAlpha(), (CAMath::Abs(prop.GetAlpha() - clAlpha) < 0.01 ? " " : " R!"), mX, mP[0], mP[1], mP[4], prop.GetQPt0(), mP[2], prop.GetSinPhi0(), "", sqrtf(mC[0]), sqrtf(mC[2]), sqrtf(mC[5]), sqrtf(mC[14]), mC[10])); // clang-format on - if (allowModification && changeDirection && !noFollowCircle && !noFollowCircle2) { + if (allowModification && false /*changeDirection*/ && !noFollowCircle && !noFollowCircle2) { if (lastRow != 255) { if (!(merger->Param().rec.tpc.disableRefitAttachment & 4)) { StoreAttachMirror(merger, lastSector, lastRow, iTrk, clAlpha, yy, xx, cluster.sector, cluster.row, inFlyDirection, prop.GetAlpha()); @@ -197,8 +178,8 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ } } else if (allowModification && lastRow != 255 && CAMath::Abs(cluster.row - lastRow) > 1) { if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { - bool dodEdx = param.dodEdxEnabled && param.rec.tpc.adddEdxSubThresholdClusters && iWay == nWays - 1 && CAMath::Abs(cluster.row - lastRow) == 2 && cluster.leg == clusters[maxN - 1].leg; - dodEdx = AttachClustersPropagate(merger, cluster.sector, lastRow, cluster.row, iTrk, cluster.leg == clusters[maxN - 1].leg, prop, inFlyDirection, GPUCA_MAX_SIN_PHI, dodEdx); + bool dodEdx = param.dodEdxEnabled && param.rec.tpc.adddEdxSubThresholdClusters && iWay == nWays - 1 && CAMath::Abs(cluster.row - lastRow) == 2; + dodEdx = AttachClustersPropagate(merger, cluster.sector, lastRow, cluster.row, iTrk, leg == 0, prop, inFlyDirection, GPUCA_MAX_SIN_PHI, dodEdx); if (dodEdx) { dEdx.fillSubThreshold(lastRow - wayDirection); if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMask) != GPUCA_GET_CONSTEXPR(param.rec.tpc, dEdxClusterRejectionFlagMaskAlt)) { @@ -244,7 +225,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ } } - if (err == 0 && changeDirection) { + if (err == 0 && false /*changeDirection*/) { const float mirrordY = prop.GetMirroredYTrack(); CADEBUG(printf(" -- MirroredY: %f --> %f", mP[0], mirrordY)); if (CAMath::Abs(yy - mP[0]) > CAMath::Abs(yy - mirrordY)) { @@ -256,7 +237,6 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ noFollowCircle = false; lastUpdateX = mX; - lastLeg = cluster.leg; lastRow = 255; N++; resetT0 = initResetT0(); @@ -270,7 +250,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ float uncorrectedY = -1e6f; if (allowModification) { - uncorrectedY = AttachClusters(merger, cluster.sector, cluster.row, iTrk, cluster.leg == clusters[maxN - 1].leg, prop); + uncorrectedY = AttachClusters(merger, cluster.sector, cluster.row, iTrk, leg == 0, prop); } const int32_t err2 = mNDF > 0 && CAMath::Abs(prop.GetSinPhi0()) >= maxSinForUpdate; @@ -334,10 +314,6 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ ConstrainSinPhi(); if (retVal == 0) // track is updated { - if (storeOuter == 1 && cluster.leg == clusters[maxN - 1].leg) { - StoreOuter(outerParam, prop, 2); - storeOuter = 255; - } noFollowCircle2 = false; lastUpdateX = mX; covYYUpd = mC[0]; @@ -352,7 +328,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_ prop.SetTrack(this, prop.GetAlpha()); } if GPUCA_RTC_CONSTEXPR (GPUCA_GET_CONSTEXPR(param.par, dodEdx)) { - if (param.dodEdxEnabled && iWay == nWays - 1 && cluster.leg == clusters[maxN - 1].leg) { // TODO: Costimize flag to remove, and option to remove double-clusters + if (param.dodEdxEnabled && iWay == nWays - 1) { // TODO: Costimize flag to remove, and option to remove double-clusters bool acc = (clusterState & param.rec.tpc.dEdxClusterRejectionFlagMask) == 0, accAlt = (clusterState & param.rec.tpc.dEdxClusterRejectionFlagMaskAlt) == 0; if (acc || accAlt) { float qtot = 0, qmax = 0, pad = 0, relTime = 0; @@ -486,7 +462,7 @@ GPUd() void GPUTPCGMTrackParam::MirrorTo(GPUTPCGMPropagator& GPUrestrict() prop, GPUd() int32_t GPUTPCGMTrackParam::MergeDoubleRowClusters(int32_t& ihit, int32_t wayDirection, GPUTPCGMMergedTrackHit* GPUrestrict() clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, const GPUTPCGMMerger* GPUrestrict() merger, GPUTPCGMPropagator& GPUrestrict() prop, float& GPUrestrict() xx, float& GPUrestrict() yy, float& GPUrestrict() zz, int32_t maxN, float clAlpha, uint8_t& GPUrestrict() clusterState, bool rejectChi2) { - if (ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector && clusters[ihit].leg == clusters[ihit + wayDirection].leg) { + if (ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector) { float maxDistY, maxDistZ; prop.GetErr2(maxDistY, maxDistZ, merger->Param(), zz, clusters[ihit].row, 0, clusters[ihit].sector, -1.f, 0.f, 0.f); // TODO: Use correct time, avgCharge maxDistY = (maxDistY + mC[0]) * 20.f; @@ -530,7 +506,7 @@ GPUd() int32_t GPUTPCGMTrackParam::MergeDoubleRowClusters(int32_t& ihit, int32_t clusterState |= clusters[ihit].state; count += clamp; } - if (!(ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector && clusters[ihit].leg == clusters[ihit + wayDirection].leg)) { + if (!(ihit + wayDirection >= 0 && ihit + wayDirection < maxN && clusters[ihit].row == clusters[ihit + wayDirection].row && clusters[ihit].sector == clusters[ihit + wayDirection].sector)) { break; } ihit += wayDirection; @@ -1071,7 +1047,7 @@ GPUd() void GPUTPCGMTrackParam::RefitTrack(GPUTPCGMMergedTrack& GPUrestrict() tr GPUTPCGMTrackParam t = track.Param(); float Alpha = track.Alpha(); CADEBUG(int32_t nTrackHitsOld = nTrackHits; float ptOld = t.QPt()); - bool ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->Param().par.earlyTpcTransform ? merger->ClustersXYZ() + track.FirstClusterRef() : nullptr, nTrackHits, NTolerated, Alpha, attempt, GPUCA_MAX_SIN_PHI, &track.OuterParam()); + bool ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->Param().par.earlyTpcTransform ? merger->ClustersXYZ() + track.FirstClusterRef() : nullptr, nTrackHits, NTolerated, Alpha, attempt, GPUCA_MAX_SIN_PHI, &track.OuterParam(), track.Leg()); CADEBUG(printf("Finished Fit Track %d\n", iTrk)); CADEBUG(printf("OUTPUT hits %d -> %d+%d = %d, QPt %f -> %f, SP %f, ok %d chi2 %f chi2ndf %f\n", nTrackHitsOld, nTrackHits, NTolerated, nTrackHits + NTolerated, ptOld, t.QPt(), t.SinPhi(), (int32_t)ok, t.Chi2(), t.Chi2() / CAMath::Max(1, nTrackHits))); @@ -1085,7 +1061,7 @@ GPUd() void GPUTPCGMTrackParam::RefitTrack(GPUTPCGMMergedTrack& GPUrestrict() tr NTolerated = 0; // Clusters not fit but tollerated for track length cut t = track.Param(); Alpha = track.Alpha(); - ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->ClustersXYZ() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, 1, GPUCA_MAX_SIN_PHI, &track.OuterParam()); + ok = t.Fit(merger, iTrk, merger->Clusters() + track.FirstClusterRef(), merger->ClustersXYZ() + track.FirstClusterRef(), nTrackHits, NTolerated, Alpha, 1, GPUCA_MAX_SIN_PHI, &track.OuterParam(), track.Leg()); } else { uint32_t nRefit = CAMath::AtomicAdd(&merger->Memory()->nRetryRefit, 1u); merger->RetryRefitIds()[nRefit] = iTrk; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h index 90ff3154a3fe9..3412388003ec6 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.h @@ -141,7 +141,7 @@ class GPUTPCGMTrackParam GPUd() bool CheckNumericalQuality(float overrideCovYY = -1.f) const; GPUd() bool CheckCov() const; - GPUd() bool Fit(GPUTPCGMMerger* merger, int32_t iTrk, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, int32_t& N, int32_t& NTolerated, float& Alpha, int32_t attempt = 0, float maxSinPhi = GPUCA_MAX_SIN_PHI, gputpcgmmergertypes::GPUTPCOuterParam* outerParam = nullptr); + GPUd() bool Fit(GPUTPCGMMerger* merger, int32_t iTrk, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, int32_t& N, int32_t& NTolerated, float& Alpha, int32_t attempt = 0, float maxSinPhi = GPUCA_MAX_SIN_PHI, gputpcgmmergertypes::GPUTPCOuterParam* outerParam = nullptr, int8_t leg = 0); GPUd() void MoveToReference(GPUTPCGMPropagator& prop, const GPUParam& param, float& alpha); GPUd() void MirrorTo(GPUTPCGMPropagator& prop, float toY, float toZ, bool inFlyDirection, const GPUParam& param, uint8_t row, uint8_t clusterState, bool mirrorParameters, int8_t sector); GPUd() int32_t MergeDoubleRowClusters(int32_t& ihit, int32_t wayDirection, GPUTPCGMMergedTrackHit* clusters, GPUTPCGMMergedTrackHitXYZ* clustersXYZ, const GPUTPCGMMerger* merger, GPUTPCGMPropagator& prop, float& xx, float& yy, float& zz, int32_t maxN, float clAlpha, uint8_t& clusterState, bool rejectChi2); diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx index 3f342c6111f04..29ccab2a765da 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx +++ b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx @@ -223,15 +223,6 @@ GPUd() int32_t GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov float tOffset; if constexpr (std::is_same_v) { count = trkX.NClusters(); - if (trkX.Looper()) { - int32_t leg = mPtrackHits[trkX.FirstClusterRef() + trkX.NClusters() - 1].leg; - for (int32_t i = trkX.NClusters() - 2; i > 0; i--) { - if (mPtrackHits[trkX.FirstClusterRef() + i].leg != leg) { - begin = i + 1; - break; - } - } - } tOffset = trkX.GetParam().GetTZOffset(); } else if constexpr (std::is_same_v) { count = trkX.getNClusters(); diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h b/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h index 431fa357e8b89..f8fa0342ee62d 100644 --- a/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h +++ b/GPU/GPUTracking/TRDTracking/GPUTRDTracker.h @@ -101,7 +101,7 @@ class GPUTRDTracker_t : public GPUProcessor { return true; } - GPUd() bool PreCheckTrackTRDCandidate(const GPUTPCGMMergedTrack& trk) const { return trk.OK() && !trk.Looper(); } + GPUd() bool PreCheckTrackTRDCandidate(const GPUTPCGMMergedTrack& trk) const { return trk.OK() && !trk.MergedLooper(); } GPUd() bool CheckTrackTRDCandidate(const TRDTRK& trk) const; GPUd() int32_t LoadTrack(const TRDTRK& trk, uint32_t tpcTrackId, bool checkTrack = true, HelperTrackAttributes* attribs = nullptr); diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 2695a732b17c7..399ed6907ad36 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -434,6 +434,16 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp // Print TPC part of track int32_t separateExtrapolatedTracksLimit = (mCfgH.separateExtrapolatedTracks ? tEXTRAPOLATEDTRACK : TRACK_TYPE_ID_LIMIT); uint32_t lastSide = -1; + int32_t prevcid = -1; + int32_t leg = 0; + if constexpr (std::is_same_v) { + if (track->PrevSegment() >= 0) { + const auto& prevtrk = mIOPtrs->mergedTracks[track->PrevSegment()]; + prevcid = mIOPtrs->mergedTrackHits[prevtrk.FirstClusterRef() + ((track->Leg() & 1) ? (prevtrk.NClusters() - 1) : 0)].num; + leg = track->Leg(); + } + } + for (int32_t k = 0; k < nClusters; k++) { if constexpr (std::is_same_v) { if (mCfgH.hideRejectedClusters && (mIOPtrs->mergedTrackHits[track->FirstClusterRef() + k].state & GPUTPCGMMergedTrackHit::flagReject)) { @@ -464,13 +474,6 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp } else { if (!drawing) { startCountInner = mVertexBuffer[iSector].size(); - if constexpr (std::is_same_v) { - if (k == 0 && track->PrevSegment() >= 0) { - const auto& prevtrk = mIOPtrs->mergedTracks[track->PrevSegment()]; - int32_t prevcid = mIOPtrs->mergedTrackHits[prevtrk.FirstClusterRef() + prevtrk.NClusters() - 1].num; - drawPointLinestrip(iSector, prevcid, tFINALTRACK, separateExtrapolatedTracksLimit); - } - } if (lastCluster != -1 && (!mCfgH.splitCETracks || lastSide == (mGlobalPos[cid].z < 0))) { int32_t lastcid; if constexpr (std::is_same_v) { @@ -479,6 +482,8 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp lastcid = &track->getCluster(mIOPtrs->outputClusRefsTPCO2, lastCluster, *mIOPtrs->clustersNative) - mIOPtrs->clustersNative->clustersLinear; } drawPointLinestrip(iSector, lastcid, tFINALTRACK, separateExtrapolatedTracksLimit); + } else if (prevcid != -1 && k == 0 && (leg & 1) == 0) { + drawPointLinestrip(iSector, prevcid, tFINALTRACK, separateExtrapolatedTracksLimit); } drawPointLinestrip(iSector, cid, tFINALTRACK, separateExtrapolatedTracksLimit); } @@ -487,6 +492,9 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp lastCluster = k; lastSide = mGlobalPos[cid].z < 0; } + if (prevcid != -1 && (leg & 1) && drawing) { + drawPointLinestrip(iSector, prevcid, tFINALTRACK, separateExtrapolatedTracksLimit); + } // Print ITS part of track if constexpr (std::is_same_v) { diff --git a/GPU/GPUTracking/qa/GPUQA.cxx b/GPU/GPUTracking/qa/GPUQA.cxx index f6ce56df6492a..35cf3e3e6e867 100644 --- a/GPU/GPUTracking/qa/GPUQA.cxx +++ b/GPU/GPUTracking/qa/GPUQA.cxx @@ -1681,7 +1681,7 @@ void GPUQA::RunQA(bool matchOnly, const std::vector* tracksEx } rowClCount += !(trackClusters[track.FirstClusterRef() + jNext].state & GPUTPCGMMergedTrackHit::flagReject); } - if (trackClusters[track.FirstClusterRef() + j].leg == trackClusters[track.FirstClusterRef() + track.NClusters() - 1].leg && rowClCount) { + if (!track.MergedLooper() && rowClCount) { nClCorrected++; } if (mcAvail && rowClCount) { From 20e0b29329e59e6c8fa638501eebf39f9b57e075 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 24 Jul 2025 12:15:36 +0200 Subject: [PATCH 11/15] GPU TPC: 16 bits are enough for nclusters --- GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h index 1ea6ab10918d5..60be206ed7e42 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergedTrack.h @@ -119,9 +119,8 @@ class GPUTPCGMMergedTrack float mAlpha; //* alpha angle uint32_t mFirstClusterRef; //* index of the first track cluster in corresponding cluster arrays int32_t mPrevSegment; //* next segment in case of looping track - // TODO: Change to 8 bit - uint32_t mNClusters; //* number of track clusters - uint32_t mNClustersFitted; //* number of clusters used in fit + uint16_t mNClusters; //* number of track clusters + uint16_t mNClustersFitted; //* number of clusters used in fit uint8_t mFlags; uint8_t mLeg; From feb430e96b14e5c8a1ebcda703393c05a03e2674 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 24 Jul 2025 13:12:21 +0200 Subject: [PATCH 12/15] GPU Display: Make 'none' frontend and backend work --- GPU/GPUTracking/Global/GPUChainTracking.cxx | 1 + .../Standalone/Benchmark/standalone.cxx | 4 +- GPU/GPUTracking/display/CMakeLists.txt | 2 + .../display/backend/GPUDisplayBackend.cxx | 4 +- .../display/backend/GPUDisplayBackend.h | 3 +- .../display/backend/GPUDisplayBackendNone.cxx | 30 ++++++++++++ .../display/backend/GPUDisplayBackendNone.h | 49 +++++++++++++++++++ .../display/frontend/GPUDisplayFrontend.cxx | 29 ++++++++++- .../display/frontend/GPUDisplayFrontend.h | 5 +- .../frontend/GPUDisplayFrontendGlfw.cxx | 11 ----- .../display/frontend/GPUDisplayFrontendGlfw.h | 1 - .../frontend/GPUDisplayFrontendGlut.cxx | 10 ---- .../display/frontend/GPUDisplayFrontendGlut.h | 1 - .../frontend/GPUDisplayFrontendNone.cxx | 19 +++++++ .../display/frontend/GPUDisplayFrontendNone.h | 7 ++- .../frontend/GPUDisplayFrontendWayland.cxx | 10 ---- .../frontend/GPUDisplayFrontendWayland.h | 1 - .../frontend/GPUDisplayFrontendWindows.cxx | 10 ---- .../frontend/GPUDisplayFrontendWindows.h | 1 - .../frontend/GPUDisplayFrontendX11.cxx | 12 +---- .../display/frontend/GPUDisplayFrontendX11.h | 6 +-- 21 files changed, 148 insertions(+), 68 deletions(-) create mode 100644 GPU/GPUTracking/display/backend/GPUDisplayBackendNone.cxx create mode 100644 GPU/GPUTracking/display/backend/GPUDisplayBackendNone.h diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx index 79e9ce6cef766..b0ea052063f20 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.cxx +++ b/GPU/GPUTracking/Global/GPUChainTracking.cxx @@ -826,6 +826,7 @@ int32_t GPUChainTracking::RunChainFinalize() if (!mDisplayRunning) { GPUInfo("Starting Event Display..."); if (mEventDisplay->StartDisplay()) { + GPUError("Error starting Event Display"); return (1); } mDisplayRunning = true; diff --git a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx index 5240b5ca47967..1b1cb510af7be 100644 --- a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx @@ -703,10 +703,10 @@ int32_t RunBenchmark(GPUReconstruction* recUse, GPUChainTracking* chainTrackingU configStandalone.noprompt = 1; } if (tmpRetVal == 3 && configStandalone.proc.ignoreNonFatalGPUErrors) { - printf("Non-FATAL GPU error occured, ignoring\n"); + printf("GPU Standalone Benchmark: Non-FATAL GPU error occured, ignoring\n"); } else if (tmpRetVal && !configStandalone.continueOnError) { if (tmpRetVal != 2) { - printf("Error occured\n"); + printf("GPU Standalone Benchmark: Error occured\n"); } return 1; } diff --git a/GPU/GPUTracking/display/CMakeLists.txt b/GPU/GPUTracking/display/CMakeLists.txt index 25b028d573bcf..32d25ee08b729 100644 --- a/GPU/GPUTracking/display/CMakeLists.txt +++ b/GPU/GPUTracking/display/CMakeLists.txt @@ -55,7 +55,9 @@ set(SRCS ../utils/qsem.cxx helpers/GPUDisplayMagneticField.cxx frontend/GPUDisplayFrontend.cxx frontend/GPUDisplayFrontendGlfw.cxx + frontend/GPUDisplayFrontendNone.cxx backend/GPUDisplayBackend.cxx + backend/GPUDisplayBackendNone.cxx backend/GPUDisplayBackendOpenGL.cxx) set(SRCS_NO_H helpers/GPUDisplayLoader.cxx diff --git a/GPU/GPUTracking/display/backend/GPUDisplayBackend.cxx b/GPU/GPUTracking/display/backend/GPUDisplayBackend.cxx index 98d2593c27950..3694ab93398cc 100644 --- a/GPU/GPUTracking/display/backend/GPUDisplayBackend.cxx +++ b/GPU/GPUTracking/display/backend/GPUDisplayBackend.cxx @@ -16,7 +16,7 @@ #include "helpers/GPUDisplayMagneticField.h" #include "GPUDisplayBackendOpenGL.h" - +#include "GPUDisplayBackendNone.h" #ifdef GPUCA_BUILD_EVENT_DISPLAY_VULKAN #include "GPUDisplayBackendVulkan.h" #endif @@ -51,6 +51,8 @@ GPUDisplayBackend* GPUDisplayBackend::getBackend(const char* type) #endif if (strcmp(type, "opengl") == 0 || strcmp(type, "auto") == 0) { return new GPUDisplayBackendOpenGL; + } else if (strcmp(type, "none") == 0) { + return new GPUDisplayBackendNone; } else { GPUError("Requested renderer not available"); } diff --git a/GPU/GPUTracking/display/backend/GPUDisplayBackend.h b/GPU/GPUTracking/display/backend/GPUDisplayBackend.h index dc56dedf587ed..546c53e1e63ff 100644 --- a/GPU/GPUTracking/display/backend/GPUDisplayBackend.h +++ b/GPU/GPUTracking/display/backend/GPUDisplayBackend.h @@ -57,7 +57,8 @@ class GPUDisplayBackend enum backendTypes { TYPE_INVALID = -1, TYPE_OPENGL = 0, - TYPE_VULKAN = 1 + TYPE_VULKAN = 1, + TYPE_NONE = 2 }; struct DrawArraysIndirectCommand { diff --git a/GPU/GPUTracking/display/backend/GPUDisplayBackendNone.cxx b/GPU/GPUTracking/display/backend/GPUDisplayBackendNone.cxx new file mode 100644 index 0000000000000..c0011265dbe52 --- /dev/null +++ b/GPU/GPUTracking/display/backend/GPUDisplayBackendNone.cxx @@ -0,0 +1,30 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file GPUDisplayBackendNone.cxx +/// \author David Rohr + +#include "GPUCommonDef.h" +#include "GPUDisplayBackendNone.h" + +using namespace o2::gpu; + +GPUDisplayBackendNone::GPUDisplayBackendNone() +{ + mBackendType = TYPE_NONE; + mBackendName = "NONE"; +} + +int32_t GPUDisplayBackendNone::InitBackendA() +{ + + return 0; +} diff --git a/GPU/GPUTracking/display/backend/GPUDisplayBackendNone.h b/GPU/GPUTracking/display/backend/GPUDisplayBackendNone.h new file mode 100644 index 0000000000000..4af69692d79c1 --- /dev/null +++ b/GPU/GPUTracking/display/backend/GPUDisplayBackendNone.h @@ -0,0 +1,49 @@ +// Copyright 2019-2020 CERN and copyright holders of ALICE O2. +// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. +// All rights not expressly granted are reserved. +// +// This software is distributed under the terms of the GNU General Public +// License v3 (GPL Version 3), copied verbatim in the file "COPYING". +// +// In applying this license CERN does not waive the privileges and immunities +// granted to it by virtue of its status as an Intergovernmental Organization +// or submit itself to any jurisdiction. + +/// \file GPUDisplayBackendNone.h +/// \author David Rohr + +#ifndef GPUDISPLAYBACKENDNONE_H +#define GPUDISPLAYBACKENDNONE_H + +#include "GPUDisplayBackend.h" + +namespace o2::gpu +{ +class GPUDisplayBackendNone : public GPUDisplayBackend +{ + public: + GPUDisplayBackendNone(); + ~GPUDisplayBackendNone() override = default; + + protected: + uint32_t DepthBits() override { return 32; }; + uint32_t drawVertices(const vboList& v, const drawType t) override { return 0; } + void ActivateColor(std::array& color) override {} + void setDepthBuffer() override {} + int32_t InitBackendA() override; + void ExitBackendA() override {} + void loadDataToGPU(size_t totalVertizes) override {} + void prepareDraw(const hmm_mat4& proj, const hmm_mat4& view, bool requestScreenshot, bool toMixBuffer, float includeMixImage) override {} + void finishDraw(bool doScreenshot, bool toMixBuffer, float includeMixImage) override {} + void finishFrame(bool doScreenshot, bool toMixBuffer, float includeMixImage) override {} + void prepareText() override {} + void finishText() override {} + void pointSizeFactor(float factor) override {} + void lineWidthFactor(float factor) override {} + void OpenGLPrint(const char* s, float x, float y, float* color, float scale) override {} + void addFontSymbol(int32_t symbol, int32_t sizex, int32_t sizey, int32_t offsetx, int32_t offsety, int32_t advance, void* data) override {} + void initializeTextDrawing() override {} +}; +} // namespace o2::gpu + +#endif diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx index 22970c3228815..6dba090a76b52 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx @@ -17,7 +17,11 @@ #ifdef _WIN32 #include "GPUDisplayFrontendWindows.h" -#elif defined(GPUCA_BUILD_EVENT_DISPLAY_X11) +#else +#include +#endif + +#ifdef GPUCA_BUILD_EVENT_DISPLAY_X11 #include "GPUDisplayFrontendX11.h" #endif #ifdef GPUCA_BUILD_EVENT_DISPLAY_GLFW @@ -29,6 +33,7 @@ #ifdef GPUCA_BUILD_EVENT_DISPLAY_WAYLAND #include "GPUDisplayFrontendWayland.h" #endif +#include "GPUDisplayFrontendNone.h" #include "GPULogging.h" #include @@ -148,7 +153,9 @@ GPUDisplayFrontend* GPUDisplayFrontend::getFrontend(const char* type) return new GPUDisplayFrontendGlut; } else #endif - { + if (strcmp(type, "none") == 0) { + return new GPUDisplayFrontendNone; + } else { GPUError("Requested frontend not available"); } return nullptr; @@ -163,3 +170,21 @@ int32_t& GPUDisplayFrontend::drawTextFontSize() { return mDisplay->drawTextFontSize(); } + +int32_t GPUDisplayFrontend::StartDisplay() +{ +#ifndef _WIN32 + static pthread_t hThread; + if (pthread_create(&hThread, nullptr, FrontendThreadWrapper, this)) { + GPUError("Coult not Create frontend Thread..."); + return (1); + } +#else + HANDLE hThread; + if ((hThread = CreateThread(nullptr, nullptr, &OpenGLWrapper, this, nullptr, nullptr)) == nullptr) { + GPUError("Coult not Create GL Thread..."); + return (1); + } +#endif + return (0); +} diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.h index 9087ec9a431f6..0abab8bb0a121 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.h @@ -40,14 +40,15 @@ class GPUDisplayFrontend : public GPUDisplayFrontendInterface TYPE_X11 = 1, TYPE_GLUT = 2, TYPE_GLFW = 3, - TYPE_WAYLAND = 4 + TYPE_WAYLAND = 4, + TYPE_NONE = 5 }; // Compile time minimum version defined in GPUDisplay.h, keep in sync! static constexpr int32_t GL_MIN_VERSION_MAJOR = 4; static constexpr int32_t GL_MIN_VERSION_MINOR = 5; - virtual int32_t StartDisplay() = 0; // Start the display. This function returns, and should spawn a thread that runs the display, and calls InitDisplay + int32_t StartDisplay(); // Start the display. This function returns, and should spawn a thread that runs the display, and calls InitDisplay void DisplayExit() override = 0; // Stop the display. Display thread should call ExitDisplay and the function returns after the thread has terminated virtual void SwitchFullscreen(bool set) = 0; // Toggle full-screen mode virtual void ToggleMaximized(bool set) = 0; // Maximize window diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.cxx index 4d80917a26215..ba22f92660fd0 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.cxx @@ -33,7 +33,6 @@ extern "C" int32_t gl3wInit(); #include #include #include -#include #ifdef GPUCA_O2_LIB #if __has_include("../src/imgui.h") @@ -417,16 +416,6 @@ void GPUDisplayFrontendGlfw::ToggleMaximized(bool set) void GPUDisplayFrontendGlfw::SetVSync(bool enable) { glfwSwapInterval(enable); } -int32_t GPUDisplayFrontendGlfw::StartDisplay() -{ - static pthread_t hThread; - if (pthread_create(&hThread, nullptr, FrontendThreadWrapper, this)) { - GPUError("Coult not Create GL Thread..."); - return (1); - } - return (0); -} - bool GPUDisplayFrontendGlfw::EnableSendKey() { #ifdef GPUCA_O2_LIB diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.h index 5276652a370a1..43dd3d65531dd 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlfw.h @@ -28,7 +28,6 @@ class GPUDisplayFrontendGlfw : public GPUDisplayFrontend GPUDisplayFrontendGlfw(); ~GPUDisplayFrontendGlfw() override = default; - int32_t StartDisplay() override; void DisplayExit() override; void SwitchFullscreen(bool set) override; void ToggleMaximized(bool set) override; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.cxx index 334a60446a4f3..1b2f2a21150c3 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.cxx @@ -309,13 +309,3 @@ void GPUDisplayFrontendGlut::SwitchFullscreen(bool set) void GPUDisplayFrontendGlut::ToggleMaximized(bool set) {} void GPUDisplayFrontendGlut::SetVSync(bool enable) {} - -int32_t GPUDisplayFrontendGlut::StartDisplay() -{ - static pthread_t hThread; - if (pthread_create(&hThread, nullptr, FrontendThreadWrapper, this)) { - GPUError("Coult not Create GL Thread..."); - return (1); - } - return (0); -} diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.h index 96f8f4af6cba5..9351349e2287d 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendGlut.h @@ -26,7 +26,6 @@ class GPUDisplayFrontendGlut : public GPUDisplayFrontend GPUDisplayFrontendGlut(); ~GPUDisplayFrontendGlut() override = default; - int32_t StartDisplay() override; void DisplayExit() override; void SwitchFullscreen(bool set) override; void ToggleMaximized(bool set) override; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.cxx index c48000bd80685..8a7eab7e00526 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.cxx @@ -13,4 +13,23 @@ /// \author David Rohr #include "GPUDisplayFrontendNone.h" +#include "GPUDisplayGUIWrapper.h" using namespace o2::gpu; + +GPUDisplayFrontendNone::GPUDisplayFrontendNone() +{ + mFrontendType = TYPE_NONE; + mFrontendName = "NONE"; +} + +int32_t GPUDisplayFrontendNone::FrontendMain() +{ + if (InitDisplay()) { + return 1; + } + do { + DrawGLScene(); + HandleSendKey(); + } while (mDisplayControl != 2); + return 0; +} diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.h index defd759ac4df6..3c7b67c35a0ce 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendNone.h @@ -21,15 +21,18 @@ namespace o2::gpu { class GPUDisplayFrontendNone : public GPUDisplayFrontend { - GPUDisplayFrontendNone() = default; + public: + GPUDisplayFrontendNone(); ~GPUDisplayFrontendNone() override = default; - int32_t StartDisplay() override { return 1; } void DisplayExit() override {} void SwitchFullscreen(bool set) override {} void ToggleMaximized(bool set) override {} void SetVSync(bool enable) override {} void OpenGLPrint(const char* s, float x, float y, float r, float g, float b, float a, bool fromBotton = true) override {} + + private: + int32_t FrontendMain() override; }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.cxx index 7a652297d89d7..5a42954c90fa7 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.cxx @@ -475,16 +475,6 @@ void GPUDisplayFrontendWayland::SetVSync(bool enable) { } -int32_t GPUDisplayFrontendWayland::StartDisplay() -{ - static pthread_t hThread; - if (pthread_create(&hThread, nullptr, FrontendThreadWrapper, this)) { - GPUError("Coult not Create frontend Thread..."); - return (1); - } - return (0); -} - void GPUDisplayFrontendWayland::getSize(int32_t& width, int32_t& height) { width = mDisplayWidth; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.h index 6dfe0a361fbb6..55676c694cfef 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWayland.h @@ -36,7 +36,6 @@ class GPUDisplayFrontendWayland : public GPUDisplayFrontend GPUDisplayFrontendWayland(); ~GPUDisplayFrontendWayland() override = default; - int32_t StartDisplay() override; void DisplayExit() override; void SwitchFullscreen(bool set) override; void ToggleMaximized(bool set) override; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.cxx index e511718e258f7..30148e0cb00ee 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.cxx @@ -375,13 +375,3 @@ void OpenGLPrint(const char* s, float x, float y, float r, float g, float b, flo void SwitchFullscreen(bool set) {} void ToggleMaximized(bool set) {} void SetVSync(bool enable) {} - -int32_t GPUDisplayFrontendWindows::StartDisplay() -{ - HANDLE hThread; - if ((hThread = CreateThread(nullptr, nullptr, &OpenGLWrapper, this, nullptr, nullptr)) == nullptr) { - GPUError("Coult not Create GL Thread..."); - return (1); - } - return (0); -} diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.h index a8534f3f0fc1f..cac5b62c4cc63 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendWindows.h @@ -25,7 +25,6 @@ class GPUDisplayFrontendWindows : public GPUDisplayFrontend GPUDisplayFrontendWindows(); ~GPUDisplayFrontendWindows() override = default; - int32_t StartDisplay() override; void DisplayExit() override; void SwitchFullscreen(bool set) override; void ToggleMaximized(bool set) override; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.cxx index 96011aa064bac..be56fc8a1e546 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.cxx @@ -23,6 +23,8 @@ #include #include +#include + #ifdef GPUCA_BUILD_EVENT_DISPLAY_VULKAN #include #include @@ -518,16 +520,6 @@ void GPUDisplayFrontendX11::SetVSync(bool enable) } } -int32_t GPUDisplayFrontendX11::StartDisplay() -{ - static pthread_t hThread; - if (pthread_create(&hThread, nullptr, FrontendThreadWrapper, this)) { - GPUError("Coult not Create frontend Thread..."); - return (1); - } - return (0); -} - void GPUDisplayFrontendX11::getSize(int32_t& width, int32_t& height) { Window root_return; diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.h b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.h index f14d05b3080bd..7ea38271f2ee9 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.h +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontendX11.h @@ -16,9 +16,10 @@ #define GPUDISPLAYFRONTENDX11_H #include "GPUDisplayFrontend.h" -#include #include -#include +#include +#include +#include #include namespace o2::gpu @@ -29,7 +30,6 @@ class GPUDisplayFrontendX11 : public GPUDisplayFrontend GPUDisplayFrontendX11(); ~GPUDisplayFrontendX11() override = default; - int32_t StartDisplay() override; void DisplayExit() override; void SwitchFullscreen(bool set) override; void ToggleMaximized(bool set) override; From 8eb87a0c4007ff174e70bd555e4b0cc6252515f5 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 24 Jul 2025 18:44:05 +0200 Subject: [PATCH 13/15] GPU Display: Skip rejected first/last clusters when drawing connected looper segments --- GPU/GPUTracking/display/render/GPUDisplayDraw.cxx | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 399ed6907ad36..359da8313274a 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -439,8 +439,14 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp if constexpr (std::is_same_v) { if (track->PrevSegment() >= 0) { const auto& prevtrk = mIOPtrs->mergedTracks[track->PrevSegment()]; - prevcid = mIOPtrs->mergedTrackHits[prevtrk.FirstClusterRef() + ((track->Leg() & 1) ? (prevtrk.NClusters() - 1) : 0)].num; leg = track->Leg(); + for (int32_t iChk = (leg & 1) ? (prevtrk.NClusters() - 1) : 0; iChk != ((leg & 1) ? -1 : (int32_t)prevtrk.NClusters()); iChk += (leg & 1) ? -1 : 1) { + const auto& hit = mIOPtrs->mergedTrackHits[prevtrk.FirstClusterRef() + iChk]; + if (!mCfgH.hideRejectedClusters || !(hit.state & GPUTPCGMMergedTrackHit::flagReject)) { + prevcid = hit.num; + break; + } + } } } From c79a241c3bbcca0358193ffbc7448a2ce1e31b6f Mon Sep 17 00:00:00 2001 From: David Rohr Date: Thu, 24 Jul 2025 22:58:52 +0200 Subject: [PATCH 14/15] GPU TPC: Fix sorting of clusters in segments of looping tracks --- GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx | 105 ++++++++++++++---- GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h | 3 - .../display/render/GPUDisplayDraw.cxx | 18 ++- 3 files changed, 92 insertions(+), 34 deletions(-) diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 81e96af47152d..0d619193fc3d2 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -1435,13 +1435,14 @@ namespace // anonymous { struct GPUTPCGMMerger_CompareClusterIds { const GPUTPCGMMerger::trackCluster* const mCmp; - GPUd() GPUTPCGMMerger_CompareClusterIds(const GPUTPCGMMerger::trackCluster* cmp) : mCmp(cmp) {} + const bool revert; + GPUd() GPUTPCGMMerger_CompareClusterIds(const GPUTPCGMMerger::trackCluster* cmp, bool r) : mCmp(cmp), revert(r) {} GPUd() bool operator()(const int16_t aa, const int16_t bb) { const GPUTPCGMMerger::trackCluster& a = mCmp[aa]; const GPUTPCGMMerger::trackCluster& b = mCmp[bb]; if (a.row != b.row) { - return (a.row > b.row); + return (a.row > b.row) ^ revert; } return GPUCA_DETERMINISTIC_CODE((a.id != b.id) ? (a.id > b.id) : (aa > bb), a.id > b.id); } @@ -1460,6 +1461,8 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread GPUTPCGMSectorTrack* trbase = nullptr; int32_t leg = 0; int32_t lastMergedSegment = -1; + bool revertSegments = false; + bool revertInSegment = false; while (true) { if (trbase && !Param().rec.tpc.dropLoopers) { int32_t jtr = trbase->NextNeighbour(); @@ -1469,7 +1472,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread trbase = nullptr; } else { trbase->SetPrevSegmentNeighbour(1000000001); - leg--; + leg += revertSegments ? 1 : -1; } } else { trbase = nullptr; @@ -1488,15 +1491,68 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread if (itr >= SectorTrackInfoLocalTotal()) { break; } - itr += nThreads * nBlocks; + revertSegments = false; + revertInSegment = false; trbase->SetPrevSegmentNeighbour(1000000000); int32_t jtr = trbase->NextNeighbour(); leg = 0; - while (jtr >= 0) { - leg++; - jtr = mSectorTrackInfos[jtr].NextNeighbour(); + if (jtr >= 0) { + int32_t lasttr = itr; + while (jtr >= 0) { // --------------- count segments --------------- + if (&mSectorTrackInfos[jtr] == trbase) { + break; // Break cyclic graph + } + lasttr = jtr; + leg++; + jtr = mSectorTrackInfos[jtr].NextNeighbour(); + } + + float mainZT = 1e9; + revertSegments = true; + for (uint32_t k = 0; k < 2; k++) { // --------------- check if first or last segment is primary --------------- + int32_t ichk = k ? lasttr : itr; + const GPUTPCGMSectorTrack* trchk = &mSectorTrackInfos[ichk]; + while (true) { + float zt = Param().par.earlyTpcTransform ? CAMath::Min(CAMath::Abs(trchk->ClusterZT0()), CAMath::Abs(trchk->ClusterZTN())) : -trchk->MinClusterZT(); // Negative time ~ smallest z, behaves the same way + if (zt < mainZT) { + if (k) { + revertSegments = false; + break; + } + mainZT = zt; + } + int32_t next = trchk->NextSegmentNeighbour(); + if (next < 0 || next == ichk) { + break; // Breaks also cycles + } + trchk = &mSectorTrackInfos[next]; + } + } + if (revertSegments) { + leg = 0; + } + + { // --------------- find longest sector track of main segment --------------- + int32_t length = 0; + int32_t ichk = revertSegments ? itr : lasttr; + const GPUTPCGMSectorTrack* trchk = &mSectorTrackInfos[ichk]; + const GPUTPCGMSectorTrack* longest = trchk; + while (true) { + if (trchk->OrigTrack()->NHits() > length) { + longest = trchk; + length = trchk->OrigTrack()->NHits(); + } + int32_t next = trchk->NextSegmentNeighbour(); + if (next < 0 || next == ichk) { + break; // Breaks also cycles + } + trchk = &mSectorTrackInfos[next]; + } + revertInSegment = (longest->ClusterZT0() < longest->ClusterZTN()) ^ (Param().par.earlyTpcTransform ? !longest->CSide() : false); + } } lastMergedSegment = -1; + itr += nThreads * nBlocks; } do { @@ -1513,7 +1569,6 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread } nHits += tr->NClusters(); - tr->SetLeg(leg); trackParts[nParts++] = tr; for (int32_t i = 0; i < 2; i++) { if (tr->ExtrapolatedTrackId(i) != -1) { @@ -1524,7 +1579,6 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread break; } trackParts[nParts] = &mSectorTrackInfos[tr->ExtrapolatedTrackId(i)]; - trackParts[nParts++]->SetLeg(leg); nHits += mSectorTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters(); } } @@ -1538,7 +1592,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread } // unpack and sort clusters - if (nParts > 1) { + if (nParts > 1 && (!revertInSegment ^ (leg & 1))) { GPUCommonAlgorithm::sort(trackParts, trackParts + nParts, [](const GPUTPCGMSectorTrack* a, const GPUTPCGMSectorTrack* b) { GPUCA_DETERMINISTIC_CODE( // clang-format off if (a->X() != b->X()) { @@ -1576,11 +1630,14 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread break; } - bool ordered = true; - for (int32_t i = 1; i < nHits; i++) { - if (trackClusters[i].row > trackClusters[i - 1].row || trackClusters[i].id == trackClusters[i - 1].id) { - ordered = false; - break; + const bool mustReverse = revertInSegment ^ (leg & 1); + bool ordered = !mustReverse; + if (ordered) { + for (int32_t i = 1; i < nHits; i++) { + if ((trackClusters[i].row > trackClusters[i - 1].row) ^ mustReverse || trackClusters[i].id == trackClusters[i - 1].id) { + ordered = false; + break; + } } } int32_t firstTrackIndex = 0; @@ -1594,7 +1651,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread clusterIndices[i] = i; } - GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIds(trackClusters)); + GPUCommonAlgorithm::sort(clusterIndices, clusterIndices + nHits, GPUTPCGMMerger_CompareClusterIds(trackClusters, mustReverse)); nTmpHits = 0; firstTrackIndex = lastTrackIndex = -1; @@ -1659,16 +1716,24 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread } GPUTPCGMMergedTrack& mergedTrack = mMergedTracks[iOutputTrack]; + GPUTPCGMTrackParam& p1 = mergedTrack.Param(); + const GPUTPCGMSectorTrack& p2 = *trackParts[firstTrackIndex]; mergedTrack.SetFlags(0); mergedTrack.SetOK(true); - mergedTrack.SetLooper(leg > 0 || lastMergedSegment >= 0); + mergedTrack.SetLeg(leg); + mergedTrack.SetLooper(leg > 0 || lastMergedSegment >= 0 || revertSegments); mergedTrack.SetNClusters(nHits); mergedTrack.SetFirstClusterRef(iMergedTrackFirstCluster); - GPUTPCGMTrackParam& p1 = mergedTrack.Param(); - const GPUTPCGMSectorTrack& p2 = *trackParts[firstTrackIndex]; mergedTrack.SetCSide(p2.CSide()); mergedTrack.SetMergedLooperConnected(leg > 0); - mergedTrack.SetPrevSegment(lastMergedSegment); + if (revertSegments) { + mergedTrack.SetPrevSegment(-1); + if (lastMergedSegment >= 0) { + mMergedTracks[lastMergedSegment].SetPrevSegment(iOutputTrack); + } + } else { + mergedTrack.SetPrevSegment(lastMergedSegment); + } lastMergedSegment = iOutputTrack; GPUTPCGMBorderTrack b; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h index 27e4a89300ca4..1de3928aac409 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.h @@ -55,7 +55,6 @@ class GPUTPCGMSectorTrack GPUd() float DzDs() const { return mParam.mDzDs; } GPUd() float QPt() const { return mParam.mQPt; } GPUd() float TZOffset() const { return mTZOffset; } - GPUd() uint8_t Leg() const { return mLeg; } GPUd() int32_t LocalTrackId() const { return mLocalTrackId; } GPUd() void SetLocalTrackId(int32_t v) { mLocalTrackId = v; } @@ -99,7 +98,6 @@ class GPUTPCGMSectorTrack GPUd() void SetNeighbor(int32_t v, int32_t i) { mNeighbour[i] = v; } GPUd() void SetPrevSegmentNeighbour(int32_t v) { mSegmentNeighbour[0] = v; } GPUd() void SetNextSegmentNeighbour(int32_t v) { mSegmentNeighbour[1] = v; } - GPUd() void SetLeg(uint8_t v) { mLeg = v; } GPUd() void CopyParamFrom(const GPUTPCGMSectorTrack& t) { @@ -136,7 +134,6 @@ class GPUTPCGMSectorTrack int32_t mLocalTrackId; // Corrected local track id in terms of GMSectorTracks array for extrapolated tracks, UNDEFINED for local tracks! int32_t mExtrapolatedTrackIds[2]; // IDs of associated extrapolated tracks uint8_t mSector; // sector of this track segment - uint8_t mLeg; // Leg of this track segment ClassDefNV(GPUTPCGMSectorTrack, 1); }; diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index 359da8313274a..7fbdae65865e2 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -376,6 +376,7 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp size_t startCountInner = mVertexBuffer[iSector].size(); bool drawing = false; + uint32_t lastSide = -1; if constexpr (std::is_same_v) { if (!mCfgH.drawTracksAndFilter && !(mCfgH.drawTPCTracks || (mCfgH.drawITSTracks && mIOPtrs->tpcLinkITS && mIOPtrs->tpcLinkITS[i] != -1) || (mCfgH.drawTRDTracks && mIOPtrs->tpcLinkTRD && mIOPtrs->tpcLinkTRD[i] != -1) || (mCfgH.drawTOFTracks && mIOPtrs->tpcLinkTOF && mIOPtrs->tpcLinkTOF[i] != -1))) { @@ -397,6 +398,7 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp drawing = true; mVertexBuffer[iSector].emplace_back(mGlobalPosTOF[cid].x, mGlobalPosTOF[cid].y * mYFactor, mCfgH.projectXY ? 0 : mGlobalPosTOF[cid].z); mGlobalPosTOF[cid].w = tTOFATTACHED; + lastSide = mGlobalPosTOF[cid].z < 0; } } @@ -410,6 +412,7 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp drawing = true; mVertexBuffer[iSector].emplace_back(mGlobalPosTRD2[cid].x, mGlobalPosTRD2[cid].y * mYFactor, mCfgH.projectXY ? 0 : mGlobalPosTRD2[cid].z); mVertexBuffer[iSector].emplace_back(mGlobalPosTRD[cid].x, mGlobalPosTRD[cid].y * mYFactor, mCfgH.projectXY ? 0 : mGlobalPosTRD[cid].z); + lastSide = mGlobalPosTRD[cid].z < 0; mGlobalPosTRD[cid].w = tTRDATTACHED; } }; @@ -433,17 +436,15 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp // Print TPC part of track int32_t separateExtrapolatedTracksLimit = (mCfgH.separateExtrapolatedTracks ? tEXTRAPOLATEDTRACK : TRACK_TYPE_ID_LIMIT); - uint32_t lastSide = -1; - int32_t prevcid = -1; - int32_t leg = 0; if constexpr (std::is_same_v) { if (track->PrevSegment() >= 0) { const auto& prevtrk = mIOPtrs->mergedTracks[track->PrevSegment()]; - leg = track->Leg(); - for (int32_t iChk = (leg & 1) ? (prevtrk.NClusters() - 1) : 0; iChk != ((leg & 1) ? -1 : (int32_t)prevtrk.NClusters()); iChk += (leg & 1) ? -1 : 1) { + for (int32_t iChk = prevtrk.NClusters() - 1; iChk >= 0; iChk--) { const auto& hit = mIOPtrs->mergedTrackHits[prevtrk.FirstClusterRef() + iChk]; if (!mCfgH.hideRejectedClusters || !(hit.state & GPUTPCGMMergedTrackHit::flagReject)) { - prevcid = hit.num; + drawPointLinestrip(iSector, hit.num, tFINALTRACK, separateExtrapolatedTracksLimit); + lastSide = mGlobalPos[hit.num].z < 0; + drawing = true; break; } } @@ -488,8 +489,6 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp lastcid = &track->getCluster(mIOPtrs->outputClusRefsTPCO2, lastCluster, *mIOPtrs->clustersNative) - mIOPtrs->clustersNative->clustersLinear; } drawPointLinestrip(iSector, lastcid, tFINALTRACK, separateExtrapolatedTracksLimit); - } else if (prevcid != -1 && k == 0 && (leg & 1) == 0) { - drawPointLinestrip(iSector, prevcid, tFINALTRACK, separateExtrapolatedTracksLimit); } drawPointLinestrip(iSector, cid, tFINALTRACK, separateExtrapolatedTracksLimit); } @@ -498,9 +497,6 @@ void GPUDisplay::DrawFinal(int32_t iSector, int32_t /*iCol*/, const GPUTPCGMProp lastCluster = k; lastSide = mGlobalPos[cid].z < 0; } - if (prevcid != -1 && (leg & 1) && drawing) { - drawPointLinestrip(iSector, prevcid, tFINALTRACK, separateExtrapolatedTracksLimit); - } // Print ITS part of track if constexpr (std::is_same_v) { From 7f38e03a7c17613d1204201d724b2d91701d4585 Mon Sep 17 00:00:00 2001 From: ALICE Action Bot Date: Thu, 24 Jul 2025 22:50:04 +0000 Subject: [PATCH 15/15] Please consider the following formatting changes --- GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx index 6dba090a76b52..d6431c9749966 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayFrontend.cxx @@ -153,7 +153,7 @@ GPUDisplayFrontend* GPUDisplayFrontend::getFrontend(const char* type) return new GPUDisplayFrontendGlut; } else #endif - if (strcmp(type, "none") == 0) { + if (strcmp(type, "none") == 0) { return new GPUDisplayFrontendNone; } else { GPUError("Requested frontend not available");