diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index 6acc7fd1dd537..a5d335931af37 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -45,7 +45,7 @@ set(SRCS SliceTracker/GPUTPCGrid.cxx SliceTracker/GPUTPCTrackletSelector.cxx SliceTracker/GPUTPCRow.cxx - SliceTracker/GPUTPCGlobalTracking.cxx + SliceTracker/GPUTPCExtrapolationTracking.cxx SliceTracker/GPUTPCCreateSliceData.cxx Merger/GPUTPCGMMerger.cxx Merger/GPUTPCGMSliceTrack.cxx diff --git a/GPU/GPUTracking/Definitions/GPUDefGPUParameters.h b/GPU/GPUTracking/Definitions/GPUDefGPUParameters.h index d246f77a50290..cebc74fcc4a5b 100644 --- a/GPU/GPUTracking/Definitions/GPUDefGPUParameters.h +++ b/GPU/GPUTracking/Definitions/GPUDefGPUParameters.h @@ -39,7 +39,7 @@ #define GPUCA_LB_GPUTPCTrackletSelector 192, 3 #define GPUCA_LB_GPUTPCNeighboursFinder 1024, 1 #define GPUCA_LB_GPUTPCNeighboursCleaner 896 - #define GPUCA_LB_GPUTPCGlobalTracking 256 + #define GPUCA_LB_GPUTPCExtrapolationTracking 256 #define GPUCA_LB_GPUTPCCFDecodeZS 64, 4 #define GPUCA_LB_GPUTPCCFDecodeZSLink GPUCA_WARP_SIZE #define GPUCA_LB_GPUTPCCFDecodeZSDenseLink GPUCA_WARP_SIZE @@ -60,7 +60,7 @@ #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step0 512 #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step2 512 #define GPUCA_LB_GPUTPCGMMergerMergeCE 512 - #define GPUCA_LB_GPUTPCGMMergerLinkGlobalTracks 256 + #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256 #define GPUCA_LB_GPUTPCGMMergerCollect 512 #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256 #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256 @@ -105,7 +105,7 @@ #define GPUCA_LB_GPUTPCTrackletSelector 256, 8 #define GPUCA_LB_GPUTPCNeighboursFinder 1024, 1 #define GPUCA_LB_GPUTPCNeighboursCleaner 896 - #define GPUCA_LB_GPUTPCGlobalTracking 256 + #define GPUCA_LB_GPUTPCExtrapolationTracking 256 #define GPUCA_LB_GPUTPCCFDecodeZS 64, 4 #define GPUCA_LB_GPUTPCCFDecodeZSLink GPUCA_WARP_SIZE #define GPUCA_LB_GPUTPCCFDecodeZSDenseLink GPUCA_WARP_SIZE @@ -126,7 +126,7 @@ #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step0 256 #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step2 256 #define GPUCA_LB_GPUTPCGMMergerMergeCE 256 - #define GPUCA_LB_GPUTPCGMMergerLinkGlobalTracks 256 + #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256 #define GPUCA_LB_GPUTPCGMMergerCollect 512 #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256 #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256 @@ -171,7 +171,7 @@ #define GPUCA_LB_GPUTPCTrackletSelector 192, 3 // best single-kernel: 128, 4 #define GPUCA_LB_GPUTPCNeighboursFinder 640, 1 // best single-kernel: 768, 1 #define GPUCA_LB_GPUTPCNeighboursCleaner 512 - #define GPUCA_LB_GPUTPCGlobalTracking 128, 4 + #define GPUCA_LB_GPUTPCExtrapolationTracking 128, 4 #define GPUCA_LB_GPUTPCCFDecodeZS 64, 10 #define GPUCA_LB_GPUTPCCFDecodeZSLink GPUCA_WARP_SIZE #define GPUCA_LB_GPUTPCCFDecodeZSDenseLink GPUCA_WARP_SIZE @@ -192,7 +192,7 @@ #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step0 192 #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step2 64, 2 #define GPUCA_LB_GPUTPCGMMergerMergeCE 256 - #define GPUCA_LB_GPUTPCGMMergerLinkGlobalTracks 256 + #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256 #define GPUCA_LB_GPUTPCGMMergerCollect 256, 2 #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256 #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256 @@ -237,7 +237,7 @@ #define GPUCA_LB_GPUTPCTrackletSelector 192, 3 #define GPUCA_LB_GPUTPCNeighboursFinder 640, 1 #define GPUCA_LB_GPUTPCNeighboursCleaner 512 - #define GPUCA_LB_GPUTPCGlobalTracking 192, 2 + #define GPUCA_LB_GPUTPCExtrapolationTracking 192, 2 #define GPUCA_LB_GPUTPCCFDecodeZS 64, 8 #define GPUCA_LB_GPUTPCCFDecodeZSLink GPUCA_WARP_SIZE #define GPUCA_LB_GPUTPCCFDecodeZSDenseLink GPUCA_WARP_SIZE @@ -258,7 +258,7 @@ #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step0 192 #define GPUCA_LB_GPUTPCGMMergerMergeBorders_step2 256 #define GPUCA_LB_GPUTPCGMMergerMergeCE 256 - #define GPUCA_LB_GPUTPCGMMergerLinkGlobalTracks 256 + #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256 #define GPUCA_LB_GPUTPCGMMergerCollect 128, 2 #define GPUCA_LB_GPUTPCGMMergerSortTracksPrepare 256 #define GPUCA_LB_GPUTPCGMMergerPrepareClusters_step0 256 @@ -312,8 +312,8 @@ #ifndef GPUCA_LB_GPUTPCNeighboursCleaner #define GPUCA_LB_GPUTPCNeighboursCleaner 256 #endif - #ifndef GPUCA_LB_GPUTPCGlobalTracking - #define GPUCA_LB_GPUTPCGlobalTracking 256 + #ifndef GPUCA_LB_GPUTPCExtrapolationTracking + #define GPUCA_LB_GPUTPCExtrapolationTracking 256 #endif #ifndef GPUCA_LB_GPUTRDTrackerKernels_gpuVersion #define GPUCA_LB_GPUTRDTrackerKernels_gpuVersion 512 @@ -414,8 +414,8 @@ #ifndef GPUCA_LB_GPUTPCGMMergerMergeCE #define GPUCA_LB_GPUTPCGMMergerMergeCE 256 #endif - #ifndef GPUCA_LB_GPUTPCGMMergerLinkGlobalTracks - #define GPUCA_LB_GPUTPCGMMergerLinkGlobalTracks 256 + #ifndef GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks + #define GPUCA_LB_GPUTPCGMMergerLinkExtrapolatedTracks 256 #endif #ifndef GPUCA_LB_GPUTPCGMMergerCollect #define GPUCA_LB_GPUTPCGMMergerCollect 256 diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index ca6f2f370300e..cd2916f5df679 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -90,8 +90,8 @@ AddOptionRTC(extraClusterErrorFactorSplitPadSharedSingleY2, float, 3.0f, "", 0, AddOptionRTC(extraClusterErrorSplitTimeSharedSingleZ2, float, 0.03f, "", 0, "Additive extra cluster error for Z2 if splittime, shared, or single set") AddOptionRTC(extraClusterErrorFactorSplitTimeSharedSingleZ2, float, 3.0f, "", 0, "Multiplicative extra cluster error for Z2 if splittime, shared, or single set") AddOptionArray(errorsCECrossing, float, 5, (0.f, 0.f, 0.f, 0.f, 0.f), "", 0, "Extra errors to add to track when crossing CE, depending on addErrorsCECrossing") // BUG: CUDA cannot yet handle AddOptionArrayRTC -AddOptionRTC(globalTrackingYRangeUpper, float, 0.85f, "", 0, "Inner portion of y-range in slice that is not used in searching for global track candidates") -AddOptionRTC(globalTrackingYRangeLower, float, 0.85f, "", 0, "Inner portion of y-range in slice that is not used in searching for global track candidates") +AddOptionRTC(extrapolationTrackingYRangeUpper, float, 0.85f, "", 0, "Inner portion of y-range in slice that is not used in searching for extrapolated track candidates") +AddOptionRTC(extrapolationTrackingYRangeLower, float, 0.85f, "", 0, "Inner portion of y-range in slice that is not used in searching for extrapolated track candidates") AddOptionRTC(trackFollowingYFactor, float, 4.f, "", 0, "Weight of y residual vs z residual in tracklet constructor") AddOptionRTC(trackMergerFactor2YS, float, 1.5f * 1.5f, "", 0, "factor2YS for track merging") AddOptionRTC(trackMergerFactor2ZT, float, 1.5f * 1.5f, "", 0, "factor2ZT for track merging") @@ -116,9 +116,9 @@ AddOptionRTC(mergerLooperSecondHorizontalQPtB5Limit, uint8_t, 2, "", 0, "Min Q/P AddOptionRTC(trackFollowingMaxRowGap, uint8_t, 4, "", 0, "Maximum number of consecutive rows without hit in track following") AddOptionRTC(trackFollowingMaxRowGapSeed, uint8_t, 2, "", 0, "Maximum number of consecutive rows without hit in track following during fit of seed") AddOptionRTC(trackFitMaxRowMissedHard, uint8_t, 10, "", 0, "Hard limit for number of missed rows in fit / propagation") -AddOptionRTC(globalTrackingRowRange, uint8_t, 45, "", 0, "Number of rows from the upped/lower limit to search for global track candidates in for") -AddOptionRTC(globalTrackingMinRows, uint8_t, 10, "", 0, "Min num of rows an additional global track must span over") -AddOptionRTC(globalTrackingMinHits, uint8_t, 8, "", 0, "Min num of hits for an additional global track") +AddOptionRTC(extrapolationTrackingRowRange, uint8_t, 45, "", 0, "Number of rows from the upped/lower limit to search for extrapolated track candidates in for") +AddOptionRTC(extrapolationTrackingMinRows, uint8_t, 10, "", 0, "Min num of rows an additional extrapolated track must span over") +AddOptionRTC(extrapolationTrackingMinHits, uint8_t, 8, "", 0, "Min num of hits for an additional extrapolated track") AddOptionRTC(noisyPadsQuickCheck, uint8_t, 1, "", 0, "Only check first fragment for noisy pads instead of all fragments (when test is enabled).") AddOptionRTC(cfQMaxCutoff, uint8_t, 3, "", 0, "Cluster Finder rejects cluster with qmax below or equal to this threshold") AddOptionRTC(cfQTotCutoff, uint8_t, 5, "", 0, "Cluster Finder rejects cluster with qtot below or equal to this threshold") @@ -134,7 +134,7 @@ AddOptionRTC(trackFitRejectMode, int8_t, 5, "", 0, "0: no limit on rejection or AddOptionRTC(rejectIFCLowRadiusCluster, uint8_t, 0, "", 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(globalTracking, int8_t, 1, "", 0, "Enable Global Tracking (prolong tracks to adjacent sectors to find short segments)") +AddOptionRTC(extrapolationTracking, int8_t, 1, "", 0, "Enable Extrapolation Tracking (prolong tracks to adjacent sectors to find short segments)") AddOptionRTC(disableRefitAttachment, uint8_t, 0, "", 0, "Bitmask to disable certain attachment steps during refit (1: attachment, 2: propagation, 4: loop following, 8: mirroring)") AddOptionRTC(rejectionStrategy, uint8_t, o2::gpu::GPUSettings::RejectionStrategyA, "", 0, "Enable rejection of TPC clusters for compression (0 = no, 1 = strategy A, 2 = strategy B)") AddOptionRTC(mergeLoopersAfterburner, uint8_t, 1, "", 0, "Run afterburner for additional looper merging") @@ -315,7 +315,7 @@ AddOption(drawInitLinks, bool, false, "", 0, "Highlight cleaned-up links") AddOption(drawSeeds, bool, false, "", 0, "Highlight seeds") AddOption(drawTracklets, bool, false, "", 0, "Highlight tracklets") AddOption(drawTracks, bool, false, "", 0, "Highlight sector tracks") -AddOption(drawGlobalTracks, bool, false, "", 0, "Highlight global sector tracks prolonged into adjacent sector") +AddOption(drawExtrapolatedTracks, bool, false, "", 0, "Highlight global sector tracks prolonged into adjacent sector") AddOption(drawFinal, bool, false, "", 0, "Highlight final tracks") AddOption(excludeClusters, int32_t, 0, "", 0, "Exclude clusters from selected draw objects from display, (2 = exclude clusters but still show tracks)") AddOption(drawSlice, int32_t, -1, "", 0, "Show individual slice") @@ -349,7 +349,7 @@ AddOption(drawTracksAndFilter, bool, false, "", 0, "Use AND filter instead of OR AddOption(propagateLoopers, bool, false, "", 0, "Enabale propagation of loopers") AddOption(clustersOnly, bool, false, "", 0, "Visualize clusters only") AddOption(clustersOnNominalRow, bool, false, "", 0, "Show clusters at nominal x of pad row for early-transformed data") -AddOption(separateGlobalTracks, bool, false, "", 0, "Draw track segments propagated to adjacent sectors separately") +AddOption(separateExtrapolatedTracks, bool, false, "", 0, "Draw track segments propagated to adjacent sectors separately") AddOption(splitCETracks, int8_t, -1, "", 0, "Split CE tracks when they cross the central electrode (-1 = for triggered data)") AddOption(markClusters, int32_t, 0, "", 0, "Mark clusters") AddOption(markFakeClusters, int32_t, 0, "", 0, "Mark fake clusters") diff --git a/GPU/GPUTracking/Global/GPUChainTracking.h b/GPU/GPUTracking/Global/GPUChainTracking.h index d827b095773b1..6e5e0b3048140 100644 --- a/GPU/GPUTracking/Global/GPUChainTracking.h +++ b/GPU/GPUTracking/Global/GPUChainTracking.h @@ -233,7 +233,7 @@ class GPUChainTracking : public GPUChain int32_t ReadEvent(uint32_t iSlice, int32_t threadId); void WriteOutput(int32_t iSlice, int32_t threadId); - int32_t GlobalTracking(uint32_t iSlice, int32_t threadId, bool synchronizeOutput = true); + int32_t ExtrapolationTracking(uint32_t iSlice, int32_t threadId, bool synchronizeOutput = true); int32_t PrepareProfile(); int32_t DoProfile(); diff --git a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx index 0831b260f881d..f28b99c0d8dd0 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingMerger.cxx @@ -161,11 +161,11 @@ int32_t GPUChainTracking::RunTPCTrackingMerger(bool synchronizeOutput) runKernel({{1, -WarpSize(), 0, deviceType, RecoStep::TPCMerging}}, MergerShadowAll.TmpCounter(), 2 * NSLICES * sizeof(*MergerShadowAll.TmpCounter())); - runKernel(GetGridAuto(0, deviceType)); + runKernel(GetGridAuto(0, deviceType)); runKernel(GetGridAuto(0, deviceType)); if (GetProcessingSettings().deterministicGPUReconstruction) { - runKernel({{1, -WarpSize(), 0, deviceType}}, 1); - runKernel({{1, -WarpSize(), 0, deviceType}}, 1); + runKernel({{1, -WarpSize(), 0, deviceType}}, 1); + runKernel({{1, -WarpSize(), 0, deviceType}}, 1); } DoDebugAndDump(RecoStep::TPCMerging, 2048, doGPU, Merger, &GPUTPCGMMerger::DumpCollected, *mDebugFile); diff --git a/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx b/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx index 174b3757d3307..b68f0797f425f 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingSliceTracker.cxx @@ -24,21 +24,13 @@ using namespace o2::gpu; -int32_t GPUChainTracking::GlobalTracking(uint32_t iSlice, int32_t threadId, bool synchronizeOutput) +int32_t GPUChainTracking::ExtrapolationTracking(uint32_t iSlice, int32_t threadId, bool synchronizeOutput) { - if (GetProcessingSettings().debugLevel >= 5) { - GPUInfo("GPU Tracker running Global Tracking for slice %u on thread %d\n", iSlice, threadId); - } - - runKernel({GetGridBlk(256, iSlice % mRec->NStreams()), {iSlice}}); + runKernel({GetGridBlk(256, iSlice % mRec->NStreams()), {iSlice}}); TransferMemoryResourceLinkToHost(RecoStep::TPCSliceTracking, processors()->tpcTrackers[iSlice].MemoryResCommon(), iSlice % mRec->NStreams()); if (synchronizeOutput) { SynchronizeStream(iSlice % mRec->NStreams()); } - - if (GetProcessingSettings().debugLevel >= 5) { - GPUInfo("GPU Tracker finished Global Tracking for slice %u on thread %d\n", iSlice, threadId); - } return (0); } @@ -64,7 +56,6 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() GPUInfo("Running TPC Slice Tracker"); } bool doGPU = GetRecoStepsGPU() & RecoStep::TPCSliceTracking; - bool doSliceDataOnGPU = processors()->tpcTrackers[0].SliceDataOnGPU(); if (!param().par.earlyTpcTransform) { for (uint32_t i = 0; i < NSLICES; i++) { processors()->tpcTrackers[i].Data().SetClusterData(nullptr, mIOPtrs.clustersNative->nClustersSector[i], mIOPtrs.clustersNative->clusterOffset[i][0]); @@ -93,7 +84,6 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() for (uint32_t iSlice = 0; iSlice < NSLICES; iSlice++) { SetupGPUProcessor(&processors()->tpcTrackers[iSlice], false); // Prepare custom allocation for 1st stack level mRec->AllocateRegisteredMemory(processors()->tpcTrackers[iSlice].MemoryResSliceScratch()); - mRec->AllocateRegisteredMemory(processors()->tpcTrackers[iSlice].MemoryResSliceInput()); } mRec->PushNonPersistentMemory(qStr2Tag("TPCSLTRK")); for (uint32_t iSlice = 0; iSlice < NSLICES; iSlice++) { @@ -173,7 +163,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() if (GetProcessingSettings().debugLevel >= 3) { GPUInfo("Creating Slice Data (Slice %d)", iSlice); } - if (doSliceDataOnGPU) { + if (doGPU) { TransferMemoryResourcesToGPU(RecoStep::TPCSliceTracking, &trk, useStream); runKernel({GetGridBlk(GPUCA_ROW_COUNT, useStream), {iSlice}, {nullptr, streamInit[useStream] ? nullptr : &mEvents->init}}); streamInit[useStream] = true; @@ -194,7 +184,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() if (GetProcessingSettings().debugLevel >= 6) { *mDebugFile << "\n\nReconstruction: Slice " << iSlice << "/" << NSLICES << std::endl; if (GetProcessingSettings().debugMask & 1) { - if (doSliceDataOnGPU) { + if (doGPU) { TransferMemoryResourcesToHost(RecoStep::TPCSliceTracking, &trk, -1, true); } trk.DumpSliceData(*mDebugFile); @@ -205,15 +195,10 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() if (GetProcessingSettings().debugLevel >= 3) { GPUInfo("Copying Slice Data to GPU and initializing temporary memory"); } - if (GetProcessingSettings().keepDisplayMemory && !doSliceDataOnGPU) { - memset((void*)trk.Data().HitWeights(), 0, trkShadow.Data().NumberOfHitsPlusAlign() * sizeof(*trkShadow.Data().HitWeights())); - } else { - runKernel(GetGridAutoStep(useStream, RecoStep::TPCSliceTracking), trkShadow.Data().HitWeights(), trkShadow.Data().NumberOfHitsPlusAlign() * sizeof(*trkShadow.Data().HitWeights())); - } + runKernel(GetGridAutoStep(useStream, RecoStep::TPCSliceTracking), trkShadow.Data().HitWeights(), trkShadow.Data().NumberOfHitsPlusAlign() * sizeof(*trkShadow.Data().HitWeights())); - // Copy Data to GPU Global Memory - if (!doSliceDataOnGPU) { - TransferMemoryResourcesToGPU(RecoStep::TPCSliceTracking, &trk, useStream); + if (!doGPU) { + TransferMemoryResourcesToGPU(RecoStep::TPCSliceTracking, &trk, useStream); // Copy Data to GPU Global Memory } if (GPUDebug("Initialization (3)", useStream)) { throw std::runtime_error("memcpy failure"); @@ -260,7 +245,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() if (!(doGPU || GetProcessingSettings().debugLevel >= 1) || GetProcessingSettings().trackletSelectorInPipeline) { runKernel({GetGridAuto(useStream), {iSlice}}); - runKernel({{1, -ThreadCount(), useStream}, {iSlice}}, 1); + runKernel({{1, -ThreadCount(), useStream}, {iSlice}}, 1); if (GetProcessingSettings().deterministicGPUReconstruction) { runKernel({GetGrid(1, 1, useStream), {iSlice}}); } @@ -317,7 +302,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() GPUInfo("Running TPC Tracklet selector (Stream %d, Slice %d to %d)", useStream, iSlice, iSlice + runSlices); } runKernel({GetGridAuto(useStream), {iSlice, runSlices}}); - runKernel({{1, -ThreadCount(), useStream}, {iSlice}}, runSlices); + runKernel({{1, -ThreadCount(), useStream}, {iSlice}}, runSlices); for (uint32_t k = iSlice; k < iSlice + runSlices; k++) { if (GetProcessingSettings().deterministicGPUReconstruction) { runKernel({GetGrid(1, 1, useStream), {k}}); @@ -337,7 +322,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() std::array transferRunning; transferRunning.fill(true); if ((GetRecoStepsOutputs() & GPUDataTypes::InOutType::TPCSectorTracks) || (doGPU && !(GetRecoStepsGPU() & RecoStep::TPCMerging))) { - if (param().rec.tpc.globalTracking) { + if (param().rec.tpc.extrapolationTracking) { mWriteOutputDone.fill(0); } @@ -386,14 +371,14 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() } mSliceSelectorReady = iSlice; - if (param().rec.tpc.globalTracking) { + if (param().rec.tpc.extrapolationTracking) { for (uint32_t tmpSlice2a = 0; tmpSlice2a <= iSlice; tmpSlice2a++) { - uint32_t tmpSlice2 = GPUTPCGlobalTracking::GlobalTrackingSliceOrder(tmpSlice2a); + uint32_t tmpSlice2 = GPUTPCExtrapolationTracking::ExtrapolationTrackingSliceOrder(tmpSlice2a); uint32_t sliceLeft, sliceRight; - GPUTPCGlobalTracking::GlobalTrackingSliceLeftRight(tmpSlice2, sliceLeft, sliceRight); + GPUTPCExtrapolationTracking::ExtrapolationTrackingSliceLeftRight(tmpSlice2, sliceLeft, sliceRight); if (tmpSlice2 <= iSlice && sliceLeft <= iSlice && sliceRight <= iSlice && mWriteOutputDone[tmpSlice2] == 0) { - GlobalTracking(tmpSlice2, 0); + ExtrapolationTracking(tmpSlice2, 0); WriteOutput(tmpSlice2, 0); mWriteOutputDone[tmpSlice2] = 1; } @@ -403,7 +388,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() } } } - if (!(GetRecoStepsOutputs() & GPUDataTypes::InOutType::TPCSectorTracks) && param().rec.tpc.globalTracking) { + if (!(GetRecoStepsOutputs() & GPUDataTypes::InOutType::TPCSectorTracks) && param().rec.tpc.extrapolationTracking) { std::vector blocking(NSLICES * mRec->NStreams()); for (int32_t i = 0; i < NSLICES; i++) { for (int32_t j = 0; j < mRec->NStreams(); j++) { @@ -411,10 +396,10 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() } } for (uint32_t iSlice = 0; iSlice < NSLICES; iSlice++) { - uint32_t tmpSlice = GPUTPCGlobalTracking::GlobalTrackingSliceOrder(iSlice); + uint32_t tmpSlice = GPUTPCExtrapolationTracking::ExtrapolationTrackingSliceOrder(iSlice); if (!((GetRecoStepsOutputs() & GPUDataTypes::InOutType::TPCSectorTracks) || (doGPU && !(GetRecoStepsGPU() & RecoStep::TPCMerging)))) { uint32_t sliceLeft, sliceRight; - GPUTPCGlobalTracking::GlobalTrackingSliceLeftRight(tmpSlice, sliceLeft, sliceRight); + GPUTPCExtrapolationTracking::ExtrapolationTrackingSliceLeftRight(tmpSlice, sliceLeft, sliceRight); if (doGPU && !blocking[tmpSlice * mRec->NStreams() + sliceLeft % mRec->NStreams()]) { StreamWaitForEvents(tmpSlice % mRec->NStreams(), &mEvents->slice[sliceLeft]); blocking[tmpSlice * mRec->NStreams() + sliceLeft % mRec->NStreams()] = true; @@ -424,7 +409,7 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() blocking[tmpSlice * mRec->NStreams() + sliceRight % mRec->NStreams()] = true; } } - GlobalTracking(tmpSlice, 0, false); + ExtrapolationTracking(tmpSlice, 0, false); } } for (uint32_t iSlice = 0; iSlice < NSLICES; iSlice++) { @@ -436,8 +421,8 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() mSliceSelectorReady = NSLICES; GPUCA_OPENMP(parallel for if(!doGPU && GetProcessingSettings().ompKernels != 1) num_threads(mRec->SetAndGetNestedLoopOmpFactor(!doGPU, NSLICES))) for (uint32_t iSlice = 0; iSlice < NSLICES; iSlice++) { - if (param().rec.tpc.globalTracking) { - GlobalTracking(iSlice, 0); + if (param().rec.tpc.extrapolationTracking) { + ExtrapolationTracking(iSlice, 0); } if (GetRecoStepsOutputs() & GPUDataTypes::InOutType::TPCSectorTracks) { WriteOutput(iSlice, 0); @@ -446,9 +431,9 @@ int32_t GPUChainTracking::RunTPCTrackingSlices_internal() mRec->SetNestedLoopOmpFactor(1); } - if (param().rec.tpc.globalTracking && GetProcessingSettings().debugLevel >= 3) { + if (param().rec.tpc.extrapolationTracking && GetProcessingSettings().debugLevel >= 3) { for (uint32_t iSlice = 0; iSlice < NSLICES; iSlice++) { - GPUInfo("Slice %d - Tracks: Local %d Global %d - Hits: Local %d Global %d", iSlice, + GPUInfo("Slice %d - Tracks: Local %d Extrapolated %d - Hits: Local %d Extrapolated %d", iSlice, processors()->tpcTrackers[iSlice].CommonMemory()->nLocalTracks, processors()->tpcTrackers[iSlice].CommonMemory()->nTracks, processors()->tpcTrackers[iSlice].CommonMemory()->nLocalTrackHits, processors()->tpcTrackers[iSlice].CommonMemory()->nTrackHits); } } diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 60dd18a254904..a0b2c7b12246a 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -502,7 +502,7 @@ GPUd() void GPUTPCGMMerger::UnpackSliceGlobal(int32_t nBlocks, int32_t nThreads, { const GPUTPCTracker& trk = GetConstantMem()->tpcTrackers[iSlice]; float alpha = Param().Alpha(iSlice); - const GPUTPCTrack* sliceTr = mMemory->firstGlobalTracks[iSlice]; + const GPUTPCTrack* sliceTr = mMemory->firstExtrapolatedTracks[iSlice]; uint32_t nLocalTracks = trk.CommonMemory()->nLocalTracks; uint32_t nTracks = *trk.NTracks(); for (uint32_t itr = nLocalTracks + iBlock * nThreads + iThread; itr < nTracks; itr += nBlocks * nThreads) { @@ -567,21 +567,21 @@ GPUd() void GPUTPCGMMerger::RefitSliceTracks(int32_t nBlocks, int32_t nThreads, track.SetNextNeighbour(-1); track.SetNextSegmentNeighbour(-1); track.SetPrevSegmentNeighbour(-1); - track.SetGlobalTrackId(0, -1); - track.SetGlobalTrackId(1, -1); + track.SetExtrapolatedTrackId(0, -1); + track.SetExtrapolatedTrackId(1, -1); uint32_t myTrack = CAMath::AtomicAdd(&mMemory->nUnpackedTracks, 1u); mTrackIDs[iSlice * mNMaxSingleSliceTracks + sliceTr->LocalTrackId()] = myTrack; mSliceTrackInfos[myTrack] = track; } } -GPUd() void GPUTPCGMMerger::LinkGlobalTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) +GPUd() void GPUTPCGMMerger::LinkExtrapolatedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread) { for (int32_t itr = SliceTrackInfoGlobalFirst(0) + iBlock * nThreads + iThread; itr < SliceTrackInfoGlobalLast(NSLICES - 1); itr += nThreads * nBlocks) { - GPUTPCGMSliceTrack& globalTrack = mSliceTrackInfos[itr]; - GPUTPCGMSliceTrack& localTrack = mSliceTrackInfos[globalTrack.LocalTrackId()]; - if (localTrack.GlobalTrackId(0) != -1 || !CAMath::AtomicCAS(&localTrack.GlobalTrackIds()[0], -1, itr)) { - localTrack.SetGlobalTrackId(1, itr); + GPUTPCGMSliceTrack& extrapolatedTrack = mSliceTrackInfos[itr]; + GPUTPCGMSliceTrack& localTrack = mSliceTrackInfos[extrapolatedTrack.LocalTrackId()]; + if (localTrack.ExtrapolatedTrackId(0) != -1 || !CAMath::AtomicCAS(&localTrack.ExtrapolatedTrackIds()[0], -1, itr)) { + localTrack.SetExtrapolatedTrackId(1, itr); } } } @@ -1521,16 +1521,16 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread tr->SetLeg(leg); trackParts[nParts++] = tr; for (int32_t i = 0; i < 2; i++) { - if (tr->GlobalTrackId(i) != -1) { + if (tr->ExtrapolatedTrackId(i) != -1) { if (nParts >= kMaxParts) { break; } - if (nHits + mSliceTrackInfos[tr->GlobalTrackId(i)].NClusters() > kMaxClusters) { + if (nHits + mSliceTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters() > kMaxClusters) { break; } - trackParts[nParts] = &mSliceTrackInfos[tr->GlobalTrackId(i)]; + trackParts[nParts] = &mSliceTrackInfos[tr->ExtrapolatedTrackId(i)]; trackParts[nParts++]->SetLeg(leg); - nHits += mSliceTrackInfos[tr->GlobalTrackId(i)].NClusters(); + nHits += mSliceTrackInfos[tr->ExtrapolatedTrackId(i)].NClusters(); } } int32_t jtr = tr->NextSegmentNeighbour(); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h index 3e4ae535fb740..7e309dcb79a9c 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.h @@ -76,7 +76,7 @@ class GPUTPCGMMerger : public GPUProcessor GPUAtomic(uint32_t) nOutputTrackClusters; GPUAtomic(uint32_t) nO2Tracks; GPUAtomic(uint32_t) nO2ClusRefs; - const GPUTPCTrack* firstGlobalTracks[NSLICES]; + const GPUTPCTrack* firstExtrapolatedTracks[NSLICES]; GPUAtomic(uint32_t) tmpCounter[2 * NSLICES]; GPUAtomic(uint32_t) nLooperMatchCandidates; }; @@ -177,7 +177,7 @@ class GPUTPCGMMerger : public GPUProcessor GPUd() void PrepareClustersForFit0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void PrepareClustersForFit1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void PrepareClustersForFit2(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); - GPUd() void LinkGlobalTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); + GPUd() void LinkExtrapolatedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void CollectMergedTracks(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void Finalize0(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); GPUd() void Finalize1(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread); diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx index c96fab2343d82..b6f11375328d0 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx @@ -137,9 +137,9 @@ GPUdii() void GPUTPCGMMergerMergeCE::Thread<0>(int32_t nBlocks, int32_t nThreads } template <> -GPUdii() void GPUTPCGMMergerLinkGlobalTracks::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger) +GPUdii() void GPUTPCGMMergerLinkExtrapolatedTracks::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger) { - merger.LinkGlobalTracks(nBlocks, nThreads, iBlock, iThread); + merger.LinkExtrapolatedTracks(nBlocks, nThreads, iBlock, iThread); } template <> diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h index e1432830117c1..dec9befa25ce2 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.h @@ -126,7 +126,7 @@ class GPUTPCGMMergerMergeCE : public GPUTPCGMMergerGeneral GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& merger); }; -class GPUTPCGMMergerLinkGlobalTracks : public GPUTPCGMMergerGeneral +class GPUTPCGMMergerLinkExtrapolatedTracks : public GPUTPCGMMergerGeneral { public: template diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSliceTrack.h b/GPU/GPUTracking/Merger/GPUTPCGMSliceTrack.h index a2179b6c66b2a..47841a616a13e 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMSliceTrack.h +++ b/GPU/GPUTracking/Merger/GPUTPCGMSliceTrack.h @@ -61,9 +61,9 @@ class GPUTPCGMSliceTrack GPUd() int32_t LocalTrackId() const { return mLocalTrackId; } GPUd() void SetLocalTrackId(int32_t v) { mLocalTrackId = v; } - GPUd() int32_t GlobalTrackId(int32_t n) const { return mGlobalTrackIds[n]; } - GPUd() void SetGlobalTrackId(int32_t n, int32_t v) { mGlobalTrackIds[n] = v; } - GPUd() int32_t* GlobalTrackIds() { return mGlobalTrackIds; } + GPUd() int32_t ExtrapolatedTrackId(int32_t n) const { return mExtrapolatedTrackIds[n]; } + GPUd() void SetExtrapolatedTrackId(int32_t n, int32_t v) { mExtrapolatedTrackIds[n] = v; } + GPUd() int32_t* ExtrapolatedTrackIds() { return mExtrapolatedTrackIds; } GPUd() float MaxClusterZT() const { return CAMath::Max(mClusterZT[0], mClusterZT[1]); } GPUd() float MinClusterZT() const { return CAMath::Min(mClusterZT[0], mClusterZT[1]); } @@ -126,19 +126,19 @@ class GPUTPCGMSliceTrack }; private: - const GPUTPCTrack* mOrigTrack; // pointer to original slice track - sliceTrackParam mParam; // Track parameters - sliceTrackParam mParam2; // Parameters at other side - float mTZOffset; // Z offset with early transform, T offset otherwise - float mAlpha; // alpha angle - float mClusterZT[2]; // Minimum maximum cluster Z / T - int32_t mNClusters; // N clusters - int32_t mNeighbour[2]; // - int32_t mSegmentNeighbour[2]; // - int32_t mLocalTrackId; // Corrected local track id in terms of GMSliceTracks array for global tracks, UNDEFINED for local tracks! - int32_t mGlobalTrackIds[2]; // IDs of associated global tracks - uint8_t mSlice; // slice of this track segment - uint8_t mLeg; // Leg of this track segment + const GPUTPCTrack* mOrigTrack; // pointer to original slice track + sliceTrackParam mParam; // Track parameters + sliceTrackParam mParam2; // Parameters at other side + float mTZOffset; // Z offset with early transform, T offset otherwise + float mAlpha; // alpha angle + float mClusterZT[2]; // Minimum maximum cluster Z / T + int32_t mNClusters; // N clusters + int32_t mNeighbour[2]; // + int32_t mSegmentNeighbour[2]; // + int32_t mLocalTrackId; // Corrected local track id in terms of GMSliceTracks array for extrapolated tracks, UNDEFINED for local tracks! + int32_t mExtrapolatedTrackIds[2]; // IDs of associated extrapolated tracks + uint8_t mSlice; // slice of this track segment + uint8_t mLeg; // Leg of this track segment ClassDefNV(GPUTPCGMSliceTrack, 1); }; diff --git a/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.cxx b/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.cxx index 9f6467923f56a..be057172a968f 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.cxx @@ -94,7 +94,7 @@ GPUdii() void GPUTPCGlobalDebugSortKernels::Thread -GPUdii() void GPUTPCGlobalDebugSortKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter) +GPUdii() void GPUTPCGlobalDebugSortKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter) { if (iThread || iBlock) { return; @@ -112,7 +112,7 @@ GPUdii() void GPUTPCGlobalDebugSortKernels::Thread -GPUdii() void GPUTPCGlobalDebugSortKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter) +GPUdii() void GPUTPCGlobalDebugSortKernels::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int8_t parameter) { if (iBlock) { return; diff --git a/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.h b/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.h index 4daee67643cfd..7c3d4246ad303 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.h +++ b/GPU/GPUTracking/Merger/GPUTPCGlobalDebugSortKernels.h @@ -29,8 +29,8 @@ class GPUTPCGlobalDebugSortKernels : public GPUKernelTemplate enum K { defaultKernel = 0, clearIds = 0, sectorTracks = 1, - globalTracks1 = 2, - globalTracks2 = 3, + extrapolatedTracks1 = 2, + extrapolatedTracks2 = 3, borderTracks = 4 }; GPUhdi() constexpr static GPUDataTypes::RecoStep GetRecoStep() { return GPUDataTypes::RecoStep::TPCMerging; } typedef GPUTPCGMMerger processorType; diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.cxx b/GPU/GPUTracking/SliceTracker/GPUTPCExtrapolationTracking.cxx similarity index 74% rename from GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.cxx rename to GPU/GPUTracking/SliceTracker/GPUTPCExtrapolationTracking.cxx index cdc72047ef0a4..1a5e99f0f52ca 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.cxx +++ b/GPU/GPUTracking/SliceTracker/GPUTPCExtrapolationTracking.cxx @@ -9,11 +9,11 @@ // granted to it by virtue of its status as an Intergovernmental Organization // or submit itself to any jurisdiction. -/// \file GPUTPCGlobalTracking.cxx +/// \file GPUTPCExtrapolationTracking.cxx /// \author David Rohr #include "GPUTPCDef.h" -#include "GPUTPCGlobalTracking.h" +#include "GPUTPCExtrapolationTracking.h" #include "GPUTPCTrackletConstructor.h" #include "GPUTPCTrackLinearisation.h" #include "GPUTPCTracker.h" @@ -22,7 +22,7 @@ using namespace o2::gpu; -GPUd() int32_t GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, const GPUTPCTracker& GPUrestrict() sliceSource, int32_t iTrack, int32_t rowIndex, float angle, int32_t direction) +GPUd() int32_t GPUTPCExtrapolationTracking::PerformExtrapolationTrackingRun(GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, const GPUTPCTracker& GPUrestrict() sliceSource, int32_t iTrack, int32_t rowIndex, float angle, int32_t direction) { /*for (int32_t j = 0;j < Tracks()[j].NHits();j++) { @@ -68,8 +68,8 @@ GPUd() int32_t GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tra } calink rowHits[GPUCA_ROW_COUNT]; - int32_t nHits = GPUTPCTrackletConstructor::GPUTPCTrackletConstructorGlobalTracking(tracker, smem, tParam, rowIndex, direction, 0, rowHits); - if (nHits >= tracker.Param().rec.tpc.globalTrackingMinHits) { + int32_t nHits = GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking(tracker, smem, tParam, rowIndex, direction, 0, rowHits); + if (nHits >= tracker.Param().rec.tpc.extrapolationTrackingMinHits) { // GPUInfo("%d hits found", nHits); uint32_t hitId = CAMath::AtomicAdd(&tracker.CommonMemory()->nTrackHits, (uint32_t)nHits); if (hitId + nHits > tracker.NMaxTrackHits()) { @@ -91,7 +91,7 @@ GPUd() int32_t GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tra if (rowHit != CALINK_INVAL && rowHit != CALINK_DEAD_CHANNEL) { // GPUInfo("New track: entry %d, row %d, hitindex %d", i, rowIndex, mTrackletRowHits[rowIndex * tracker.CommonMemory()->nTracklets]); tracker.TrackHits()[hitId + i].Set(rowIndex, rowHit); - // if (i == 0) tParam.TransportToX(Row(rowIndex).X(), Param().bzCLight(), GPUCA_MAX_SIN_PHI); //Use transport with new linearisation, we have changed the track in between - NOT needed, fitting will always start at outer end of global track! + // if (i == 0) tParam.TransportToX(Row(rowIndex).X(), Param().bzCLight(), GPUCA_MAX_SIN_PHI); //Use transport with new linearisation, we have changed the track in between - NOT needed, fitting will always start at outer end of the extrapolated track! i++; } rowIndex++; @@ -115,42 +115,42 @@ GPUd() int32_t GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tra track.SetLocalTrackId((sliceSource.ISlice() << 24) | sliceSource.Tracks()[iTrack].LocalTrackId()); } - return (nHits >= tracker.Param().rec.tpc.globalTrackingMinHits); + return (nHits >= tracker.Param().rec.tpc.extrapolationTrackingMinHits); } -GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, const GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, GPUTPCTracker& GPUrestrict() sliceTarget, bool right) +GPUd() void GPUTPCExtrapolationTracking::PerformExtrapolationTracking(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, const GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, GPUTPCTracker& GPUrestrict() sliceTarget, bool right) { for (int32_t i = iBlock * nThreads + iThread; i < tracker.CommonMemory()->nLocalTracks; i += nThreads * nBlocks) { { const int32_t tmpHit = tracker.Tracks()[i].FirstHitID(); - if (tracker.TrackHits()[tmpHit].RowIndex() >= tracker.Param().rec.tpc.globalTrackingMinRows && tracker.TrackHits()[tmpHit].RowIndex() < tracker.Param().rec.tpc.globalTrackingRowRange) { + if (tracker.TrackHits()[tmpHit].RowIndex() >= tracker.Param().rec.tpc.extrapolationTrackingMinRows && tracker.TrackHits()[tmpHit].RowIndex() < tracker.Param().rec.tpc.extrapolationTrackingRowRange) { int32_t rowIndex = tracker.TrackHits()[tmpHit].RowIndex(); const GPUTPCRow& GPUrestrict() row = tracker.Row(rowIndex); float Y = (float)tracker.Data().HitDataY(row, tracker.TrackHits()[tmpHit].HitIndex()) * row.HstepY() + row.Grid().YMin(); - if (!right && Y < -row.MaxY() * tracker.Param().rec.tpc.globalTrackingYRangeLower) { + if (!right && Y < -row.MaxY() * tracker.Param().rec.tpc.extrapolationTrackingYRangeLower) { // GPUInfo("Track %d, lower row %d, left border (%f of %f)", i, mTrackHits[tmpHit].RowIndex(), Y, -row.MaxY()); - PerformGlobalTrackingRun(sliceTarget, smem, tracker, i, rowIndex, -tracker.Param().par.dAlpha, -1); + PerformExtrapolationTrackingRun(sliceTarget, smem, tracker, i, rowIndex, -tracker.Param().par.dAlpha, -1); } - if (right && Y > row.MaxY() * tracker.Param().rec.tpc.globalTrackingYRangeLower) { + if (right && Y > row.MaxY() * tracker.Param().rec.tpc.extrapolationTrackingYRangeLower) { // GPUInfo("Track %d, lower row %d, right border (%f of %f)", i, mTrackHits[tmpHit].RowIndex(), Y, row.MaxY()); - PerformGlobalTrackingRun(sliceTarget, smem, tracker, i, rowIndex, tracker.Param().par.dAlpha, -1); + PerformExtrapolationTrackingRun(sliceTarget, smem, tracker, i, rowIndex, tracker.Param().par.dAlpha, -1); } } } { const int32_t tmpHit = tracker.Tracks()[i].FirstHitID() + tracker.Tracks()[i].NHits() - 1; - if (tracker.TrackHits()[tmpHit].RowIndex() < GPUCA_ROW_COUNT - tracker.Param().rec.tpc.globalTrackingMinRows && tracker.TrackHits()[tmpHit].RowIndex() >= GPUCA_ROW_COUNT - tracker.Param().rec.tpc.globalTrackingRowRange) { + if (tracker.TrackHits()[tmpHit].RowIndex() < GPUCA_ROW_COUNT - tracker.Param().rec.tpc.extrapolationTrackingMinRows && tracker.TrackHits()[tmpHit].RowIndex() >= GPUCA_ROW_COUNT - tracker.Param().rec.tpc.extrapolationTrackingRowRange) { int32_t rowIndex = tracker.TrackHits()[tmpHit].RowIndex(); const GPUTPCRow& GPUrestrict() row = tracker.Row(rowIndex); float Y = (float)tracker.Data().HitDataY(row, tracker.TrackHits()[tmpHit].HitIndex()) * row.HstepY() + row.Grid().YMin(); - if (!right && Y < -row.MaxY() * tracker.Param().rec.tpc.globalTrackingYRangeUpper) { + if (!right && Y < -row.MaxY() * tracker.Param().rec.tpc.extrapolationTrackingYRangeUpper) { // GPUInfo("Track %d, upper row %d, left border (%f of %f)", i, mTrackHits[tmpHit].RowIndex(), Y, -row.MaxY()); - PerformGlobalTrackingRun(sliceTarget, smem, tracker, i, rowIndex, -tracker.Param().par.dAlpha, 1); + PerformExtrapolationTrackingRun(sliceTarget, smem, tracker, i, rowIndex, -tracker.Param().par.dAlpha, 1); } - if (right && Y > row.MaxY() * tracker.Param().rec.tpc.globalTrackingYRangeUpper) { + if (right && Y > row.MaxY() * tracker.Param().rec.tpc.extrapolationTrackingYRangeUpper) { // GPUInfo("Track %d, upper row %d, right border (%f of %f)", i, mTrackHits[tmpHit].RowIndex(), Y, row.MaxY()); - PerformGlobalTrackingRun(sliceTarget, smem, tracker, i, rowIndex, tracker.Param().par.dAlpha, 1); + PerformExtrapolationTrackingRun(sliceTarget, smem, tracker, i, rowIndex, tracker.Param().par.dAlpha, 1); } } } @@ -158,7 +158,7 @@ GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int32_t nBlocks, int32_t } template <> -GPUdii() void GPUTPCGlobalTracking::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() tracker) +GPUdii() void GPUTPCExtrapolationTracking::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() tracker) { CA_SHARED_CACHE(&smem.mRows[0], tracker.SliceDataRows(), GPUCA_ROW_COUNT * sizeof(GPUTPCRow)); GPUbarrier(); @@ -173,11 +173,11 @@ GPUdii() void GPUTPCGlobalTracking::Thread<0>(int32_t nBlocks, int32_t nThreads, sliceLeft += GPUDataTypes::NSLICES / 2; sliceRight += GPUDataTypes::NSLICES / 2; } - PerformGlobalTracking(nBlocks, nThreads, iBlock, iThread, tracker.GetConstantMem()->tpcTrackers[sliceLeft], smem, tracker, true); - PerformGlobalTracking(nBlocks, nThreads, iBlock, iThread, tracker.GetConstantMem()->tpcTrackers[sliceRight], smem, tracker, false); + PerformExtrapolationTracking(nBlocks, nThreads, iBlock, iThread, tracker.GetConstantMem()->tpcTrackers[sliceLeft], smem, tracker, true); + PerformExtrapolationTracking(nBlocks, nThreads, iBlock, iThread, tracker.GetConstantMem()->tpcTrackers[sliceRight], smem, tracker, false); } -GPUd() int32_t GPUTPCGlobalTracking::GlobalTrackingSliceOrder(int32_t iSlice) +GPUd() int32_t GPUTPCExtrapolationTracking::ExtrapolationTrackingSliceOrder(int32_t iSlice) { iSlice++; if (iSlice == GPUDataTypes::NSLICES / 2) { @@ -189,7 +189,7 @@ GPUd() int32_t GPUTPCGlobalTracking::GlobalTrackingSliceOrder(int32_t iSlice) return iSlice; } -GPUd() void GPUTPCGlobalTracking::GlobalTrackingSliceLeftRight(uint32_t iSlice, uint32_t& left, uint32_t& right) +GPUd() void GPUTPCExtrapolationTracking::ExtrapolationTrackingSliceLeftRight(uint32_t iSlice, uint32_t& left, uint32_t& right) { left = (iSlice + (GPUDataTypes::NSLICES / 2 - 1)) % (GPUDataTypes::NSLICES / 2); right = (iSlice + 1) % (GPUDataTypes::NSLICES / 2); @@ -200,7 +200,7 @@ GPUd() void GPUTPCGlobalTracking::GlobalTrackingSliceLeftRight(uint32_t iSlice, } template <> -GPUdii() void GPUTPCGlobalTrackingCopyNumbers::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() tracker, int32_t n) +GPUdii() void GPUTPCExtrapolationTrackingCopyNumbers::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() tracker, int32_t n) { for (int32_t i = get_global_id(0); i < n; i += get_global_size(0)) { GPUconstantref() GPUTPCTracker& GPUrestrict() trk = (&tracker)[i]; diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.h b/GPU/GPUTracking/SliceTracker/GPUTPCExtrapolationTracking.h similarity index 68% rename from GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.h rename to GPU/GPUTracking/SliceTracker/GPUTPCExtrapolationTracking.h index c3f765f42cec5..cd6533a3439ed 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCExtrapolationTracking.h @@ -9,11 +9,11 @@ // granted to it by virtue of its status as an Intergovernmental Organization // or submit itself to any jurisdiction. -/// \file GPUTPCGlobalTracking.h +/// \file GPUTPCExtrapolationTracking.h /// \author David Rohr -#ifndef GPUTPCGLOBALTRACKING_H -#define GPUTPCGLOBALTRACKING_H +#ifndef GPUTPCEXTRAPOLATIONTRACKING_H +#define GPUTPCEXTRAPOLATIONTRACKING_H #include "GPUGeneralKernels.h" #include "GPUConstantMem.h" @@ -24,7 +24,7 @@ namespace gpu { class GPUTPCTracker; -class GPUTPCGlobalTracking : public GPUKernelTemplate +class GPUTPCExtrapolationTracking : public GPUKernelTemplate { public: struct GPUSharedMemory { @@ -40,15 +40,15 @@ class GPUTPCGlobalTracking : public GPUKernelTemplate template GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& tracker); - GPUd() static int32_t GlobalTrackingSliceOrder(int32_t iSlice); - GPUd() static void GlobalTrackingSliceLeftRight(uint32_t iSlice, uint32_t& left, uint32_t& right); + GPUd() static int32_t ExtrapolationTrackingSliceOrder(int32_t iSlice); + GPUd() static void ExtrapolationTrackingSliceLeftRight(uint32_t iSlice, uint32_t& left, uint32_t& right); private: - GPUd() static int32_t PerformGlobalTrackingRun(GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, const GPUTPCTracker& sliceSource, int32_t iTrack, int32_t rowIndex, float angle, int32_t direction); - GPUd() static void PerformGlobalTracking(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, const GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, GPUTPCTracker& sliceTarget, bool right); + GPUd() static int32_t PerformExtrapolationTrackingRun(GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, const GPUTPCTracker& sliceSource, int32_t iTrack, int32_t rowIndex, float angle, int32_t direction); + GPUd() static void PerformExtrapolationTracking(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, const GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& smem, GPUTPCTracker& sliceTarget, bool right); }; -class GPUTPCGlobalTrackingCopyNumbers : public GPUKernelTemplate +class GPUTPCExtrapolationTrackingCopyNumbers : public GPUKernelTemplate { public: typedef GPUconstantref() GPUTPCTracker processorType; diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx index e02cba2144920..8a727dc2da930 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx +++ b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx @@ -58,20 +58,6 @@ void GPUTPCSliceData::SetMaxData() mNumberOfHitsPlusAlign = GPUProcessor::nextMultipleOf<(kVectorAlignment > GPUCA_ROWALIGNMENT ? kVectorAlignment : GPUCA_ROWALIGNMENT) / sizeof(int32_t)>(hitMemCount); } -void* GPUTPCSliceData::SetPointersInput(void* mem, bool idsOnGPU, bool sliceDataOnGPU) -{ - if (sliceDataOnGPU) { - return mem; - } - const int32_t firstHitInBinSize = GetGridSize(mNumberOfHits, GPUCA_ROW_COUNT) + GPUCA_ROW_COUNT * GPUCA_ROWALIGNMENT / sizeof(int32_t); - GPUProcessor::computePointerWithAlignment(mem, mHitData, mNumberOfHitsPlusAlign); - GPUProcessor::computePointerWithAlignment(mem, mFirstHitInBin, firstHitInBinSize); - if (idsOnGPU) { - mem = SetPointersClusterIds(mem, false); // Hijack the allocation from SetPointersClusterIds - } - return mem; -} - void* GPUTPCSliceData::SetPointersLinks(void* mem) { GPUProcessor::computePointerWithAlignment(mem, mLinkUpData, mNumberOfHitsPlusAlign); @@ -85,10 +71,13 @@ void* GPUTPCSliceData::SetPointersWeights(void* mem) return mem; } -void* GPUTPCSliceData::SetPointersScratch(void* mem, bool idsOnGPU, bool sliceDataOnGPU) +void* GPUTPCSliceData::SetPointersScratch(void* mem, bool idsOnGPU) { - if (sliceDataOnGPU) { - mem = SetPointersInput(mem, idsOnGPU, false); + const int32_t firstHitInBinSize = GetGridSize(mNumberOfHits, GPUCA_ROW_COUNT) + GPUCA_ROW_COUNT * GPUCA_ROWALIGNMENT / sizeof(int32_t); + GPUProcessor::computePointerWithAlignment(mem, mHitData, mNumberOfHitsPlusAlign); + GPUProcessor::computePointerWithAlignment(mem, mFirstHitInBin, firstHitInBinSize); + if (idsOnGPU) { + mem = SetPointersClusterIds(mem, false); // Hijack the allocation from SetPointersClusterIds } return mem; } diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.h b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.h index 72e9f9d2c19d5..200a123b9bb83 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCSliceData.h @@ -38,8 +38,7 @@ class GPUTPCSliceData void InitializeRows(const GPUParam& p); void SetMaxData(); void SetClusterData(const GPUTPCClusterData* data, int32_t nClusters, int32_t clusterIdOffset); - void* SetPointersInput(void* mem, bool idsOnGPU, bool sliceDataOnGPU); - void* SetPointersScratch(void* mem, bool idsOnGPU, bool sliceDataOnGPU); + void* SetPointersScratch(void* mem, bool idsOnGPU); void* SetPointersLinks(void* mem); void* SetPointersWeights(void* mem); void* SetPointersClusterIds(void* mem, bool idsOnGPU); diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTrack.h b/GPU/GPUTracking/SliceTracker/GPUTPCTrack.h index 18418bc031d7e..fcf9d1149c588 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTrack.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTrack.h @@ -62,7 +62,7 @@ class GPUTPCTrack private: int32_t mFirstHitID; // index of the first track cell in the track->cell pointer array int32_t mNHits; // number of track cells - int32_t mLocalTrackId; // Id of local track this global track belongs to, index of this track itself if it is a local track + int32_t mLocalTrackId; // Id of local track this extrapolated track belongs to, index of this track itself if it is a local track GPUTPCBaseTrackParam mParam; // track parameters private: diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTracker.cxx b/GPU/GPUTracking/SliceTracker/GPUTPCTracker.cxx index d5a941b333c6e..df0c7813fa0db 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTracker.cxx +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTracker.cxx @@ -40,7 +40,7 @@ using namespace o2::tpc; #if !defined(GPUCA_GPUCODE) GPUTPCTracker::GPUTPCTracker() - : GPUProcessor(), mLinkTmpMemory(nullptr), mISlice(-1), mData(), mNMaxStartHits(0), mNMaxRowStartHits(0), mNMaxTracklets(0), mNMaxRowHits(0), mNMaxTracks(0), mNMaxTrackHits(0), mMemoryResLinks(-1), mMemoryResScratchHost(-1), mMemoryResCommon(-1), mMemoryResTracklets(-1), mMemoryResOutput(-1), mMemoryResSliceScratch(-1), mMemoryResSliceInput(-1), mRowStartHitCountOffset(nullptr), mTrackletTmpStartHits(nullptr), mGPUTrackletTemp(nullptr), mGPUParametersConst(), mCommonMem(nullptr), mTrackletStartHits(nullptr), mTracklets(nullptr), mTrackletRowHits(nullptr), mTracks(nullptr), mTrackHits(nullptr), mOutput(nullptr), mOutputMemory(nullptr) + : GPUProcessor(), mLinkTmpMemory(nullptr), mISlice(-1), mData(), mNMaxStartHits(0), mNMaxRowStartHits(0), mNMaxTracklets(0), mNMaxRowHits(0), mNMaxTracks(0), mNMaxTrackHits(0), mMemoryResLinks(-1), mMemoryResScratchHost(-1), mMemoryResCommon(-1), mMemoryResTracklets(-1), mMemoryResOutput(-1), mMemoryResSliceScratch(-1), mRowStartHitCountOffset(nullptr), mTrackletTmpStartHits(nullptr), mGPUTrackletTemp(nullptr), mGPUParametersConst(), mCommonMem(nullptr), mTrackletStartHits(nullptr), mTracklets(nullptr), mTrackletRowHits(nullptr), mTracks(nullptr), mTrackHits(nullptr), mOutput(nullptr), mOutputMemory(nullptr) { } @@ -62,15 +62,9 @@ void GPUTPCTracker::InitializeProcessor() SetupCommonMemory(); } -bool GPUTPCTracker::SliceDataOnGPU() -{ - return (mRec->GetRecoStepsGPU() & GPUDataTypes::RecoStep::TPCSliceTracking) && (mRec->GetRecoStepsGPU() & GPUDataTypes::RecoStep::TPCConversion) && (mRec->GetConstantMem().ioPtrs.clustersNative || mRec->GetConstantMem().ioPtrs.tpcZS || mRec->GetConstantMem().ioPtrs.tpcPackedDigits); -} - -void* GPUTPCTracker::SetPointersDataInput(void* mem) { return mData.SetPointersInput(mem, mRec->GetRecoStepsGPU() & GPUDataTypes::RecoStep::TPCMerging, SliceDataOnGPU()); } void* GPUTPCTracker::SetPointersDataLinks(void* mem) { return mData.SetPointersLinks(mem); } void* GPUTPCTracker::SetPointersDataWeights(void* mem) { return mData.SetPointersWeights(mem); } -void* GPUTPCTracker::SetPointersDataScratch(void* mem) { return mData.SetPointersScratch(mem, mRec->GetRecoStepsGPU() & GPUDataTypes::RecoStep::TPCMerging, SliceDataOnGPU()); } +void* GPUTPCTracker::SetPointersDataScratch(void* mem) { return mData.SetPointersScratch(mem, mRec->GetRecoStepsGPU() & GPUDataTypes::RecoStep::TPCMerging); } void* GPUTPCTracker::SetPointersDataRows(void* mem) { return mData.SetPointersRows(mem); } void* GPUTPCTracker::SetPointersScratch(void* mem) @@ -108,7 +102,6 @@ void GPUTPCTracker::RegisterMemoryAllocation() GPUMemoryReuse reLinks{reuseCondition, GPUMemoryReuse::REUSE_1TO1, GPUMemoryReuse::TrackerDataLinks, (uint16_t)(mISlice % mRec->GetProcessingSettings().nStreams)}; mMemoryResLinks = mRec->RegisterMemoryAllocation(this, &GPUTPCTracker::SetPointersDataLinks, GPUMemoryResource::MEMORY_SCRATCH | GPUMemoryResource::MEMORY_STACK, "TPCSliceLinks", reLinks); mMemoryResSliceScratch = mRec->RegisterMemoryAllocation(this, &GPUTPCTracker::SetPointersDataScratch, GPUMemoryResource::MEMORY_SCRATCH | GPUMemoryResource::MEMORY_STACK | GPUMemoryResource::MEMORY_CUSTOM, "TPCSliceScratch"); - mMemoryResSliceInput = mRec->RegisterMemoryAllocation(this, &GPUTPCTracker::SetPointersDataInput, GPUMemoryResource::MEMORY_INPUT | GPUMemoryResource::MEMORY_STACK | GPUMemoryResource::MEMORY_CUSTOM, "TPCSliceInput"); GPUMemoryReuse reWeights{reuseCondition, GPUMemoryReuse::REUSE_1TO1, GPUMemoryReuse::TrackerDataWeights, (uint16_t)(mISlice % mRec->GetProcessingSettings().nStreams)}; mRec->RegisterMemoryAllocation(this, &GPUTPCTracker::SetPointersDataWeights, GPUMemoryResource::MEMORY_SCRATCH | GPUMemoryResource::MEMORY_STACK, "TPCSliceWeights", reWeights); GPUMemoryReuse reScratch{reuseCondition, GPUMemoryReuse::REUSE_1TO1, GPUMemoryReuse::TrackerScratch, (uint16_t)(mISlice % mRec->GetProcessingSettings().nStreams)}; diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTracker.h b/GPU/GPUTracking/SliceTracker/GPUTPCTracker.h index 10259c80ac80c..c5d4d40a2bef8 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTracker.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTracker.h @@ -79,7 +79,7 @@ class GPUTPCTracker : public GPUProcessor GPUAtomic(uint32_t) nTracklets; // number of tracklets GPUAtomic(uint32_t) nRowHits; // number of tracklet hits GPUAtomic(uint32_t) nTracks; // number of reconstructed tracks - int32_t nLocalTracks; // number of reconstructed tracks before global tracking + int32_t nLocalTracks; // number of reconstructed tracks before extrapolation tracking GPUAtomic(uint32_t) nTrackHits; // number of track hits int32_t nLocalTrackHits; // see above StructGPUParameters gpuParameters; // GPU parameters @@ -114,8 +114,6 @@ class GPUTPCTracker : public GPUProcessor } void SetupCommonMemory(); - bool SliceDataOnGPU(); - void* SetPointersDataInput(void* mem); void* SetPointersDataLinks(void* mem); void* SetPointersDataWeights(void* mem); void* SetPointersDataScratch(void* mem); @@ -133,7 +131,6 @@ class GPUTPCTracker : public GPUProcessor int16_t MemoryResTracklets() const { return mMemoryResTracklets; } int16_t MemoryResOutput() const { return mMemoryResOutput; } int16_t MemoryResSliceScratch() const { return mMemoryResSliceScratch; } - int16_t MemoryResSliceInput() const { return mMemoryResSliceInput; } void SetMaxData(const GPUTrackingInOutPointers& io); void UpdateMaxData(); @@ -257,7 +254,6 @@ class GPUTPCTracker : public GPUProcessor int16_t mMemoryResTracklets; int16_t mMemoryResOutput; int16_t mMemoryResSliceScratch; - int16_t mMemoryResSliceInput; // GPU Temp Arrays GPUglobalref() int32_t* mRowStartHitCountOffset; // Offset, length and new offset of start hits in row diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.cxx index 8e8c82393d659..04833375ad6df 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.cxx +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.cxx @@ -21,7 +21,7 @@ #include "GPUTPCTracker.h" #include "GPUTPCTracklet.h" #include "GPUTPCTrackletConstructor.h" -#include "GPUTPCGlobalTracking.h" +#include "GPUTPCExtrapolationTracking.h" #include "CorrectionMapsHelper.h" #include "CalibdEdxContainer.h" #include "GPUParam.inc" @@ -565,7 +565,7 @@ GPUd() int32_t GPUTPCTrackletConstructor::FetchTracklet(GPUconstantref() GPUTPCT #endif // GPUCA_GPUCODE template <> // FIXME: GPUgeneric() needed to make the clang spirv output link correctly -GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorGlobalTracking(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCGlobalTracking::GPUSharedMemory& sMem, GPUTPCTrackParam& GPUrestrict() tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits) +GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, GPUTPCTrackParam& GPUrestrict() tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits) { GPUTPCThreadMemory rMem; rMem.mISH = iTracklet; diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.h b/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.h index 88a2d9c94d305..9af1eeb0ae7b2 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.h @@ -96,7 +96,7 @@ class GPUTPCTrackletConstructor #endif // GPUCA_GPUCODE template - GPUd() static int32_t GPUTPCTrackletConstructorGlobalTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, GPUTPCTrackParam& tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits); + GPUd() static int32_t GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, GPUTPCTrackParam& tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits); typedef GPUconstantref() GPUTPCTracker processorType; GPUhdi() constexpr static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::TPCSliceTracking; } diff --git a/GPU/GPUTracking/Standalone/tools/dump.C b/GPU/GPUTracking/Standalone/tools/dump.C deleted file mode 100644 index f5213f40a8a95..0000000000000 --- a/GPU/GPUTracking/Standalone/tools/dump.C +++ /dev/null @@ -1,7 +0,0 @@ -void dump() -{ - AliHLTSystem* pHLT = AliHLTPluginBase::GetInstance(); - AliHLTConfiguration overrideClusterTransformation("TPC-ClusterTransformation", "TPCClusterTransformation", "TPC-HWCFDecoder", "-use-orig-transform -do-mc"); - AliHLTConfiguration dumper("Dumper", "GPUDump", "TPC-ClusterTransformation TRD-tracklet-reader", ""); - AliHLTConfiguration overrideTracker("TPC-TR", "TPCCATracker", "TPC-ClusterTransformation Dumper", "-GlobalTracking -SearchWindowDZDR 2.5"); -} diff --git a/GPU/GPUTracking/display/GPUDisplay.cxx b/GPU/GPUTracking/display/GPUDisplay.cxx index 4c770b32ee66a..918011b85ea04 100644 --- a/GPU/GPUTracking/display/GPUDisplay.cxx +++ b/GPU/GPUTracking/display/GPUDisplay.cxx @@ -516,17 +516,17 @@ void GPUDisplay::DrawGLScene_drawCommands() LOOP_SLICE LOOP_COLLISION_COL(drawVertices(mGlDLPoints[iSlice][tSLICETRACK][iCol], GPUDisplayBackend::POINTS)); skip2:; - if (mCfgL.drawGlobalTracks) { + if (mCfgL.drawExtrapolatedTracks) { if (mCfgL.excludeClusters) { goto skip3; } if (mCfgL.colorClusters) { - SetColorGlobalTracks(); + SetColorExtrapolatedTracks(); } } else { SetColorClusters(); } - LOOP_SLICE LOOP_COLLISION_COL(drawVertices(mGlDLPoints[iSlice][tGLOBALTRACK][iCol], GPUDisplayBackend::POINTS)); + LOOP_SLICE LOOP_COLLISION_COL(drawVertices(mGlDLPoints[iSlice][tEXTRAPOLATEDTRACK][iCol], GPUDisplayBackend::POINTS)); SetColorClusters(); if (mCfgL.drawFinal && mCfgL.propagateTracks < 2) { @@ -564,9 +564,9 @@ void GPUDisplay::DrawGLScene_drawCommands() SetColorTracks(); LOOP_SLICE drawVertices(mGlDLLines[iSlice][tSLICETRACK], GPUDisplayBackend::LINE_STRIP); } - if (mCfgL.drawGlobalTracks) { - SetColorGlobalTracks(); - LOOP_SLICE drawVertices(mGlDLLines[iSlice][tGLOBALTRACK], GPUDisplayBackend::LINE_STRIP); + if (mCfgL.drawExtrapolatedTracks) { + SetColorExtrapolatedTracks(); + LOOP_SLICE drawVertices(mGlDLLines[iSlice][tEXTRAPOLATEDTRACK], GPUDisplayBackend::LINE_STRIP); } } if (mCfgL.drawFinal) { @@ -670,7 +670,7 @@ void GPUDisplay::DrawGLScene_internal(float animateTime, bool renderToMixBuffer) snprintf(info, 1024, "FPS: %6.2f (Slice: %d, 1:Clusters %d, 2:Prelinks %d, 3:Links %d, 4:Seeds %d, 5:Tracklets %d, 6:Tracks %d, 7:GTracks %d, 8:Merger %d) (%d frames, %d draw calls) " "(X %1.2f Y %1.2f Z %1.2f / R %1.2f Phi %1.1f Theta %1.1f) / Yaw %1.1f Pitch %1.1f Roll %1.1f)", - fps, mCfgL.drawSlice, mCfgL.drawClusters, mCfgL.drawInitLinks, mCfgL.drawLinks, mCfgL.drawSeeds, mCfgL.drawTracklets, mCfgL.drawTracks, mCfgL.drawGlobalTracks, mCfgL.drawFinal, mFramesDone, mNDrawCalls, mXYZ[0], mXYZ[1], mXYZ[2], mRPhiTheta[0], mRPhiTheta[1] * 180 / CAMath::Pi(), + fps, mCfgL.drawSlice, mCfgL.drawClusters, mCfgL.drawInitLinks, mCfgL.drawLinks, mCfgL.drawSeeds, mCfgL.drawTracklets, mCfgL.drawTracks, mCfgL.drawExtrapolatedTracks, mCfgL.drawFinal, mFramesDone, mNDrawCalls, mXYZ[0], mXYZ[1], mXYZ[2], mRPhiTheta[0], mRPhiTheta[1] * 180 / CAMath::Pi(), mRPhiTheta[2] * 180 / CAMath::Pi(), mAngle[1] * 180 / CAMath::Pi(), mAngle[0] * 180 / CAMath::Pi(), mAngle[2] * 180 / CAMath::Pi()); if (fpstime > 1.) { if (mPrintInfoText & 2) { diff --git a/GPU/GPUTracking/display/GPUDisplay.h b/GPU/GPUTracking/display/GPUDisplay.h index e7836461e4fd9..d6a65f212ecf3 100644 --- a/GPU/GPUTracking/display/GPUDisplay.h +++ b/GPU/GPUTracking/display/GPUDisplay.h @@ -96,7 +96,7 @@ class GPUDisplay : public GPUDisplayInterface tSEED = 3, tTRACKLET = 4, tSLICETRACK = 5, - tGLOBALTRACK = 6, + tEXTRAPOLATEDTRACK = 6, tFINALTRACK = 7, tMARKED = 8, tTRDCLUSTER = 9, @@ -188,7 +188,7 @@ class GPUDisplay : public GPUDisplayInterface void SetColorSeeds(); void SetColorTracklets(); void SetColorTracks(); - void SetColorGlobalTracks(); + void SetColorExtrapolatedTracks(); void SetColorFinal(); void SetColorGrid(); void SetColorGridTRD(); diff --git a/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx b/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx index 431240e93b732..acf5566489f49 100644 --- a/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx +++ b/GPU/GPUTracking/display/frontend/GPUDisplayKeys.cxx @@ -60,7 +60,7 @@ const char* HelpText[] = { "[SHIFT] Slow Zoom / Move / Rotate", "[ALT] / [CTRL] / [ENTER] Focus camera on origin / orient y-axis upwards (combine with [SHIFT] to lock) / Cycle through modes", "[RCTRL] / [RALT] Rotate model instead of camera / rotate TPC around beamline", - "[1] ... [8] / [N] Enable display of clusters, preseeds, seeds, starthits, tracklets, tracks, global tracks, merged tracks / Show assigned clusters in colors", + "[1] ... [8] / [N] Enable display of clusters, preseeds, seeds, starthits, tracklets, tracks, extrapolated tracks, merged tracks / Show assigned clusters in colors", "[F1] / [F2] / [F3] / [F4] Enable / disable drawing of TPC / TRD / TOF / ITS", "[SHIFT] + [F1] to [F4] Enable / disable track detector filter", "[SHIFT] + [F12] Switch track detector filter between AND and OR mode" @@ -164,11 +164,11 @@ void GPUDisplay::HandleKey(uint8_t key) mPrintInfoText &= 3; SetInfo("Info text display - console: %s, onscreen %s", (mPrintInfoText & 2) ? "enabled" : "disabled", (mPrintInfoText & 1) ? "enabled" : "disabled"); } else if (key == 'j') { - if (mCfgH.separateGlobalTracks) { + if (mCfgH.separateExtrapolatedTracks) { mCfgH.splitCETracks ^= 1; } - mCfgH.separateGlobalTracks ^= 1; - SetInfo("Seperated display of tracks propagated to adjacent sectors %s / of CE tracks %s", mCfgH.separateGlobalTracks ? "enabled" : "disabled", mCfgH.splitCETracks ? "enabled" : "disabled"); + mCfgH.separateExtrapolatedTracks ^= 1; + SetInfo("Seperated display of tracks propagated to adjacent sectors %s / of CE tracks %s", mCfgH.separateExtrapolatedTracks ? "enabled" : "disabled", mCfgH.splitCETracks ? "enabled" : "disabled"); } else if (key == 'c') { if (mCfgH.markClusters == 0) { mCfgH.markClusters = 1; @@ -310,7 +310,7 @@ void GPUDisplay::HandleKey(uint8_t key) } else if (key == '6') { mCfgL.drawTracks ^= 1; } else if (key == '7') { - mCfgL.drawGlobalTracks ^= 1; + mCfgL.drawExtrapolatedTracks ^= 1; } else if (key == '8') { mCfgL.drawFinal ^= 1; } else if (key == mFrontend->KEY_F1) { diff --git a/GPU/GPUTracking/display/helpers/GPUDisplayColors.inc b/GPU/GPUTracking/display/helpers/GPUDisplayColors.inc index c10e0d3a55876..3716a07536e20 100644 --- a/GPU/GPUTracking/display/helpers/GPUDisplayColors.inc +++ b/GPU/GPUTracking/display/helpers/GPUDisplayColors.inc @@ -114,7 +114,7 @@ inline void GPUDisplay::SetColorTracks() } ActivateColor(); } -inline void GPUDisplay::SetColorGlobalTracks() +inline void GPUDisplay::SetColorExtrapolatedTracks() { if (mCfgL.invertColors) { mDrawColor = {0.8, 0.2, 0, 1.f}; diff --git a/GPU/GPUTracking/display/helpers/GPUDisplayHelpers.cxx b/GPU/GPUTracking/display/helpers/GPUDisplayHelpers.cxx index 866d4a59aab82..764f659d07e64 100644 --- a/GPU/GPUTracking/display/helpers/GPUDisplayHelpers.cxx +++ b/GPU/GPUTracking/display/helpers/GPUDisplayHelpers.cxx @@ -52,7 +52,7 @@ void GPUDisplay::disableUnsupportedOptions() mCfgH.markFakeClusters = 0; } if (!mChain) { - mCfgL.excludeClusters = mCfgL.drawInitLinks = mCfgL.drawLinks = mCfgL.drawSeeds = mCfgL.drawTracklets = mCfgL.drawTracks = mCfgL.drawGlobalTracks = 0; + mCfgL.excludeClusters = mCfgL.drawInitLinks = mCfgL.drawLinks = mCfgL.drawSeeds = mCfgL.drawTracklets = mCfgL.drawTracks = mCfgL.drawExtrapolatedTracks = 0; } if (mConfig.showTPCTracksFromO2Format && mParam->par.earlyTpcTransform) { throw std::runtime_error("Cannot run GPU display with early Transform when input is O2 tracks"); diff --git a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx index b1685fc61fc2c..5d4628cf0eb3f 100644 --- a/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx +++ b/GPU/GPUTracking/display/render/GPUDisplayDraw.cxx @@ -414,7 +414,7 @@ void GPUDisplay::DrawFinal(int32_t iSlice, int32_t /*iCol*/, GPUTPCGMPropagator* } // Print TPC part of track - int32_t separateGlobalTracksLimit = (mCfgH.separateGlobalTracks ? tGLOBALTRACK : TRACK_TYPE_ID_LIMIT); + int32_t separateExtrapolatedTracksLimit = (mCfgH.separateExtrapolatedTracks ? tEXTRAPOLATEDTRACK : TRACK_TYPE_ID_LIMIT); uint32_t lastSide = -1; for (int32_t k = 0; k < nClusters; k++) { if constexpr (std::is_same_v) { @@ -435,10 +435,10 @@ void GPUDisplay::DrawFinal(int32_t iSlice, int32_t /*iCol*/, GPUTPCGMPropagator* drawing = false; lastCluster = -1; } else { - drawPointLinestrip(iSlice, cid, tFINALTRACK, separateGlobalTracksLimit); + drawPointLinestrip(iSlice, cid, tFINALTRACK, separateExtrapolatedTracksLimit); } } - if (w == separateGlobalTracksLimit) { + if (w == separateExtrapolatedTracksLimit) { if (drawing) { insertVertexList(vBuf[0], startCountInner, mVertexBuffer[iSlice].size()); } @@ -453,9 +453,9 @@ void GPUDisplay::DrawFinal(int32_t iSlice, int32_t /*iCol*/, GPUTPCGMPropagator* } else { lastcid = &track->getCluster(mIOPtrs->outputClusRefsTPCO2, lastCluster, *mIOPtrs->clustersNative) - mIOPtrs->clustersNative->clustersLinear; } - drawPointLinestrip(iSlice, lastcid, tFINALTRACK, separateGlobalTracksLimit); + drawPointLinestrip(iSlice, lastcid, tFINALTRACK, separateExtrapolatedTracksLimit); } - drawPointLinestrip(iSlice, cid, tFINALTRACK, separateGlobalTracksLimit); + drawPointLinestrip(iSlice, cid, tFINALTRACK, separateExtrapolatedTracksLimit); } drawing = true; } @@ -812,7 +812,7 @@ size_t GPUDisplay::DrawGLScene_updateVertexList() GPUCA_OPENMP(for) for (int32_t iSlice = 0; iSlice < NSLICES; iSlice++) { const GPUTPCTracker& tracker = sliceTracker(iSlice); - mGlDLLines[iSlice][tGLOBALTRACK] = DrawTracks(tracker, 1); + mGlDLLines[iSlice][tEXTRAPOLATEDTRACK] = DrawTracks(tracker, 1); } GPUCA_OPENMP(barrier) } diff --git a/GPU/GPUTracking/kernels.cmake b/GPU/GPUTracking/kernels.cmake index 57f0cce4989f3..4085bebee08c4 100644 --- a/GPU/GPUTracking/kernels.cmake +++ b/GPU/GPUTracking/kernels.cmake @@ -36,16 +36,16 @@ o2_gpu_add_kernel("GPUTPCTrackletConstructor, allSlices" "= TPCTRAC o2_gpu_add_kernel("GPUTPCTrackletSelector" "= TPCTRACKER" LB both) o2_gpu_add_kernel("GPUMemClean16" "GPUGeneralKernels" NO "simple, REG, (GPUCA_THREAD_COUNT, 1)" void* ptr "uint64_t" size) o2_gpu_add_kernel("GPUitoa" "GPUGeneralKernels" NO "simple, REG, (GPUCA_THREAD_COUNT, 1)" int32_t* ptr "uint64_t" size) -o2_gpu_add_kernel("GPUTPCGlobalTrackingCopyNumbers" "GPUTPCGlobalTracking TPCTRACKER" NO single int32_t n) -o2_gpu_add_kernel("GPUTPCGlobalTracking" "= TPCTRACKER TPCTRACKLETCONS" LB single) +o2_gpu_add_kernel("GPUTPCExtrapolationTrackingCopyNumbers" "GPUTPCExtrapolationTracking TPCTRACKER" NO single int32_t n) +o2_gpu_add_kernel("GPUTPCExtrapolationTracking" "= TPCTRACKER TPCTRACKLETCONS" LB single) o2_gpu_add_kernel("GPUTPCCreateSliceData" "= TPCTRACKER TPCSLICEDATA" LB single) o2_gpu_add_kernel("GPUTPCSectorDebugSortKernels, hitData" "= TPCTRACKER" NO single) o2_gpu_add_kernel("GPUTPCSectorDebugSortKernels, startHits" "= TPCTRACKER" NO single) o2_gpu_add_kernel("GPUTPCSectorDebugSortKernels, sliceTracks" "= TPCTRACKER" NO single) o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, clearIds" "= TPCMERGER" NO single int8_t parameter) o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, sectorTracks" "= TPCMERGER" NO single int8_t parameter) -o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, globalTracks1" "= TPCMERGER" NO single int8_t parameter) -o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, globalTracks2" "= TPCMERGER" NO single int8_t parameter) +o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, extrapolatedTracks1" "= TPCMERGER" NO single int8_t parameter) +o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, extrapolatedTracks2" "= TPCMERGER" NO single int8_t parameter) o2_gpu_add_kernel("GPUTPCGlobalDebugSortKernels, borderTracks" "= TPCMERGER" NO single int8_t parameter) o2_gpu_add_kernel("GPUTPCCreateOccupancyMap, fill" "= TPCOCCUPANCY" LB simple GPUTPCClusterOccupancyMapBin* map) o2_gpu_add_kernel("GPUTPCCreateOccupancyMap, fold" "= TPCOCCUPANCY" LB simple GPUTPCClusterOccupancyMapBin* map "uint32_t*" output) @@ -68,7 +68,7 @@ o2_gpu_add_kernel("GPUTPCGMMergerMergeBorders, step1" "GPUTPCGMM o2_gpu_add_kernel("GPUTPCGMMergerMergeBorders, step2" "GPUTPCGMMergerGPU TPCMERGER" LB simple int32_t iSlice int8_t withinSlice int8_t mergeMode) o2_gpu_add_kernel("GPUTPCGMMergerMergeBorders, variant" "GPUTPCGMMergerGPU TPCMERGER" NO simple gputpcgmmergertypes::GPUTPCGMBorderRange* range int32_t N int32_t cmpMax) o2_gpu_add_kernel("GPUTPCGMMergerMergeCE" "GPUTPCGMMergerGPU TPCMERGER" LB simple) -o2_gpu_add_kernel("GPUTPCGMMergerLinkGlobalTracks" "GPUTPCGMMergerGPU TPCMERGER" LB simple) +o2_gpu_add_kernel("GPUTPCGMMergerLinkExtrapolatedTracks" "GPUTPCGMMergerGPU TPCMERGER" LB simple) o2_gpu_add_kernel("GPUTPCGMMergerCollect" "GPUTPCGMMergerGPU TPCMERGER" LB simple) o2_gpu_add_kernel("GPUTPCGMMergerSortTracks" "GPUTPCGMMergerGPU TPCMERGER" NO simple) o2_gpu_add_kernel("GPUTPCGMMergerSortTracksQPt" "GPUTPCGMMergerGPU TPCMERGER" NO simple) diff --git a/cmake/O2RootMacroExclusionList.cmake b/cmake/O2RootMacroExclusionList.cmake index 4b87da5b4e42e..d5596ccc424f4 100644 --- a/cmake/O2RootMacroExclusionList.cmake +++ b/cmake/O2RootMacroExclusionList.cmake @@ -42,7 +42,6 @@ list(APPEND O2_ROOT_MACRO_EXCLUSION_LIST GPU/GPUTracking/Merger/macros/fitPolynomialFieldIts.C # Needs AliRoot AliMagF GPU/GPUTracking/Merger/macros/fitPolynomialFieldTpc.C # Needs AliRoot AliMagF GPU/GPUTracking/Merger/macros/fitPolynomialFieldTrd.C # Needs AliRoot AliMagF - GPU/GPUTracking/Standalone/tools/dump.C # Needs AliRoot ALiHLTSystem GPU/GPUTracking/Standalone/tools/dumpTRDClusterMatrices.C # Needs AliRoot AliCDBManager, AliGeomManager and AliTRDgeometry GPU/GPUTracking/TRDTracking/macros/checkDbgOutput.C # Needs AliRoot TStatToolkit GPU/TPCFastTransformation/devtools/loadlibs.C # Special macro