ITS: new CPU + GPU seeding vertexer - #15733
Conversation
Adds a seeding vertexer that runs as a prepended tracker pass (diamond trackleting -> cells -> lines -> parallel seeding), on both the CPU and GPU traits, replacing the per-ROF CPU vertexer for the seeding step.
5c4ab24 to
06b0b02
Compare
| // MC only | ||
| template <int NLayers> | ||
| void TimeFrameGPU<NLayers>::downloadMemberOffsetsDevice() | ||
| { | ||
| const int nPeaks = mPeakOffsetsHost.empty() ? 0 : mPeakOffsetsHost.back(); | ||
| mMemberOffsetsHost.resize(nPeaks + 1); | ||
| if (nPeaks) { | ||
| GPUChkErrS(cudaMemcpyAsync(mMemberOffsetsHost.data(), mMemberOffsetsDevice, (nPeaks + 1) * sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[0].get())); | ||
| GPUChkErrS(cudaStreamSynchronize(mGpuStreams[0].get())); | ||
| } | ||
| mNMembers = (nPeaks && !mMemberOffsetsHost.empty()) ? mMemberOffsetsHost.back() : 0; | ||
| } | ||
|
|
||
| template <int NLayers> | ||
| void TimeFrameGPU<NLayers>::downloadMemberLinesDevice() | ||
| { | ||
| mMemberLinesHost.resize(mNMembers); | ||
| if (mNMembers) { | ||
| GPUChkErrS(cudaMemcpyAsync(mMemberLinesHost.data(), mMemberLinesDevice, mNMembers * sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[0].get())); | ||
| GPUChkErrS(cudaStreamSynchronize(mGpuStreams[0].get())); | ||
| } | ||
| } |
There was a problem hiding this comment.
In general, I would not support MC labels on the device but simply provide the ability to download them and rebuilt the MC on demand
There was a problem hiding this comment.
Do we need to duplicate the already existing ClusterLines math, can we not have only one of these classes?
| void prepareClusters(const TrackingParameters& trkParam, const int maxLayers) override | ||
| { | ||
| if (maxLayers < NLayers) { // only if former seeding vertexer is run | ||
| TimeFrame<NLayers>::prepareClusters(trkParam, maxLayers); | ||
| } | ||
| } | ||
| void allocateClusterSortStorage(const TrackingParameters& trkParam, const int maxLayers) override | ||
| { | ||
| if (maxLayers < NLayers) { // only if former seeding vertexer is run | ||
| TimeFrame<NLayers>::allocateClusterSortStorage(trkParam, maxLayers); | ||
| } | ||
| } |
| std::bitset<NLayers> mUnsortedClustersUploaded{}; | ||
| std::bitset<NLayers> mTrackingFrameInfoUploaded{}; |
There was a problem hiding this comment.
These are not used anywhere?
| std::vector<o2::MCCompLabel> mLineLabelFlatHost; | ||
| int mNMembers{0}; | ||
| Vertex* mDiamondDevice{nullptr}; | ||
| bool mPersistentTablesLoaded{false}; |
There was a problem hiding this comment.
Do we actually need this flag, can we not just add to the configuration the equivalent step?
| bool useParallelSeeding{false}; // use the GPU-oriented parallel seeding (computeVertices) | ||
| bool incrementalSeeding{false}; // CPU-only fast path for computeVertices |
There was a problem hiding this comment.
Why is this needed, we know by construction if we are running on CPU or GPU why do we need a configKeyVal?
| float lineMinPt = -1.f; | ||
| float fineZWindow = -1.f; | ||
| int fineMinDensity = 8; | ||
| float fineMaxDrift = -1.f; | ||
| float goodLineChi2Cut = 5.f; | ||
| float goodLinePtCut = 0.5f; | ||
| float goodContributorsSignificance = -1.f; | ||
| float duplicateZScale = -1.f; |
There was a problem hiding this comment.
In general all the config key values should be set to the Pb-Pb default values, we only really override them for the pp.
| inline bool timeCompatible(const LineTime& a, const LineTime& b) | ||
| { | ||
| return o2::gpu::GPUCommonMath::Abs(a.tc - b.tc) <= (a.th + b.th); | ||
| } | ||
|
|
There was a problem hiding this comment.
if we anyways have LineTime class why not make this a function of that class?
Also If you were to use the already existing TimeEstBC class then we not even need that.
| void computeDensities(const bounded_vector<float>& Z, | ||
| const bounded_vector<LineTime>& T, | ||
| const float zWindow, | ||
| const bool incremental, | ||
| bounded_vector<LineWindow>& win, | ||
| bounded_vector<int>& density) | ||
| { | ||
| const int m = static_cast<int>(Z.size()); | ||
| int lo{0}, hi{0}; | ||
| for (int k = 0; k < m; ++k) { | ||
| const float zk = Z[k]; | ||
| if (incremental) { | ||
| while (lo < m && Z[lo] < zk - zWindow) { | ||
| ++lo; | ||
| } | ||
| while (hi < m && Z[hi] <= zk + zWindow) { | ||
| ++hi; | ||
| } | ||
| } else { | ||
| lo = static_cast<int>(std::lower_bound(Z.begin(), Z.end(), zk - zWindow) - Z.begin()); | ||
| hi = static_cast<int>(std::upper_bound(Z.begin(), Z.end(), zk + zWindow) - Z.begin()); | ||
| } | ||
| win[k] = LineWindow{lo, hi}; | ||
| int count = 0; | ||
| for (int j = lo; j < hi; ++j) { | ||
| count += timeCompatible(T[k], T[j]); | ||
| } | ||
| density[k] = count; | ||
| } | ||
| } | ||
|
|
||
| void markLeftmostMaxima(const bounded_vector<int>& density, | ||
| const bounded_vector<float>& Z, | ||
| const bounded_vector<LineWindow>& win, | ||
| const bool incremental, | ||
| std::pmr::memory_resource* pool, | ||
| bounded_vector<uint8_t>& isMax) | ||
| { | ||
| const int m = static_cast<int>(density.size()); | ||
| if (incremental) { | ||
| bounded_vector<int> maxDeque(m, pool); | ||
| int front{0}, back{0}, next{0}; | ||
| for (int k = 0; k < m; ++k) { | ||
| while (next < win[k].hi) { | ||
| while (back > front && density[maxDeque[back - 1]] < density[next]) { | ||
| --back; | ||
| } | ||
| maxDeque[back++] = next++; | ||
| } | ||
| while (front < back && maxDeque[front] < win[k].lo) { | ||
| ++front; | ||
| } | ||
| isMax[k] = maxDeque[front] == k; // the window always contains k itself, so the deque is non-empty | ||
| } | ||
| return; | ||
| } | ||
| for (int k = 0; k < m; ++k) { | ||
| uint8_t best = 1; | ||
| for (int j = win[k].lo; j < win[k].hi; ++j) { | ||
| if (j == k) { | ||
| continue; | ||
| } | ||
| if (density[j] > density[k] || (density[j] == density[k] && Z[j] < Z[k])) { | ||
| best = 0; | ||
| break; | ||
| } | ||
| } | ||
| isMax[k] = best; | ||
| } | ||
| } | ||
| } // namespace |
There was a problem hiding this comment.
These functions are only called on the host, so incremental is always true?
| float deltaPhi{o2::gpu::CAMath::Abs(currentTracklet.phi - nextTracklet.phi)}; | ||
| if (deltaPhi > o2::constants::math::PI) { | ||
| deltaPhi = o2::constants::math::TwoPI - deltaPhi; | ||
| } | ||
| if (deltaPhi > cellDeltaPhiCut) { | ||
| continue; | ||
| } |
There was a problem hiding this comment.
in MathUtils.h exist isPhiDifferenceBelow
|
@shahor02 This algorithm should be tested, but I have some preliminary results. On 50 simulated PbPb TFs at 50kHz:
On a PbPb TF with embedded pp collisions:PbPb:
pp:
pp track efficiency:
|
|
Thanks, looks good! Note that there are conflicts in the PR. |
Yes absolutely |
Sorry for replying late: next week we are considering to move the ITS WP2 meeting from Wednesday 2 september to Thursday 3 September, starting at 10 am. |
|
@fprino I will be on vacation from Thursday onwards. |
|
Some things from an initial test: [PROD][fschlepp@epn000 pvtest]$ python3 time.py gabriele/reco.log
Processing:
Original count: 925
Trimmed count: 649
Discarded: 138 low + 138 high
Total: 336405.15 ms (336.41 s)
Average: 518.34 ms
Std deviation: 194.22 ms
Vertex seeding:
Original count: 925
Trimmed count: 649
Discarded: 138 low + 138 high
Total: 60625.25 ms (60.63 s)
Average: 93.41 ms
Std deviation: 13.94 ms
Tracking:
Original count: 919
Trimmed count: 645
Discarded: 137 low + 137 high
Total: 231575.46 ms (231.58 s)
Average: 359.03 ms
Std deviation: 86.99 ms
[PROD][fschlepp@epn000 pvtest]$ python3 time.py dev/reco.log
Processing:
Original count: 925
Trimmed count: 649
Discarded: 138 low + 138 high
Total: 511572.46 ms (511.57 s)
Average: 788.25 ms
Std deviation: 239.93 ms
Vertex seeding:
Original count: 925
Trimmed count: 649
Discarded: 138 low + 138 high
Total: 222618.46 ms (222.62 s)
Average: 343.02 ms
Std deviation: 89.72 ms
Tracking:
Original count: 924
Trimmed count: 648
Discarded: 138 low + 138 high
Total: 240011.05 ms (240.01 s)
Average: 370.39 ms
Std deviation: 70.58 msnote this timing from MI50 EPNs. the new vertexing seems to be 3.7 times faster which in total means a speedup of 1.5 in total processing time (-34.2%). 23 kHz Pb-Pb. black:dev, red:this PR, blue:this PR but relaxing ITSVertexerParam.clusterContributorsCut from 3 to 2: Note that there is currently no second iteration for UPC flagging. |
|
OK, I saw https://github.com/AliceO2Group/AliceO2/pull/15733/changes#diff-27fc8aaeb7769a496f838182d30a759ba63b3163bc7f458dd9bbbc895d3c7211R258-R260, but then the ITSVertexerParam are irrelevant in that mode, and there is no ITS trackfinding UPC iteration? |
There is no track finding since no ROFs are flagged by a vertex carrying the UPC flag in the UPC mask, all ROFs are automatically masked in UPC iteration |
Hi, sorry today I am travelling so I won't be very responsive. I didn't put yet the UPC iteration in the PR, I will test one in the next days. suppressLowMultDebris=13 slipped when copying the parameters, but in any case it is not used by this algorithm by default, thanks for noticing. |






Adds an optional seeding vertexer that runs as a prepended tracker pass (diamond trackleting -> cells -> lines -> parallel seeding), on both the CPU and GPU traits.
Additionally, unsorted clusters are loaded to GPU and then sorted on GPU.
To set it:
ITSCATrackerParam.seedingVertexIteration=1;
ITSVertexerParam.useParallelSeeding=1;
Tuned parameters for Pb--Pb:
ITSCATrackerParam.seedingVertexIteration=1;ITSVertexerParam.useParallelSeeding=1;ITSVertexerParam.useTruthSeeding=0;ITSCATrackerParam.diamondTrackletingPVres=1.5102214480735774;ITSCATrackerParam.diamondTrackletingNSigmaCut=2.976946815605669;ITSCATrackerParam.diamondTrackletingCellDeltaTanLambdaSigma=0.005566295924057422;ITSCATrackerParam.diamondCellTanLambdaNSigma=2.5;ITSCATrackerParam.diamondTrackletingCellDeltaPhiMinPt=0.12;ITSCATrackerParam.cellLineSharedClusterCut=1;ITSVertexerParam.clusterCut=0.07;ITSVertexerParam.pairCut=0.025684202919354766;ITSVertexerParam.phiCut=0.008875191490279988;ITSVertexerParam.clusterContributorsCut=3;ITSVertexerParam.lineMinPt=0.10;ITSVertexerParam.nSigmaCut=0.0016;ITSVertexerParam.goodLineChi2Cut=9.814679066181695;ITSVertexerParam.goodLinePtCut=0.12;ITSVertexerParam.goodContributorsSignificance=0.070;ITSVertexerParam.suppressLowMultDebris=13;ITSVertexerParam.fineZWindow=0.010;ITSVertexerParam.fineMinDensity=8;ITSVertexerParam.fineMaxDrift=0.005;ITSVertexerParam.duplicateZScale=0.7;ITSVertexerParam.duplicateZCut=0.1388966993405415;
Tuned parameters for pp:
ITSCATrackerParam.seedingVertexIteration=1;ITSVertexerParam.useParallelSeeding=1;ITSVertexerParam.useTruthSeeding=0;ITSCATrackerParam.diamondTrackletingPVres=3.74344532796542;ITSCATrackerParam.diamondTrackletingNSigmaCut=5.352263543811316;ITSCATrackerParam.diamondTrackletingCellDeltaTanLambdaSigma=0.002563403956702525;ITSCATrackerParam.diamondCellTanLambdaNSigma=3.257473162361808;ITSCATrackerParam.diamondTrackletingCellDeltaPhiMinPt=0.09407482537938365;ITSCATrackerParam.cellLineSharedClusterCut=3;ITSVertexerParam.clusterCut=0.05557266167329535;ITSVertexerParam.pairCut=0.05176031785438137;ITSVertexerParam.phiCut=0.017639175978851104;ITSVertexerParam.clusterContributorsCut=2;ITSVertexerParam.lineMinPt=0.1;ITSVertexerParam.nSigmaCut=0.0016;ITSVertexerParam.goodLineChi2Cut=9.814679066181695;ITSVertexerParam.goodLinePtCut=0.12;ITSVertexerParam.goodContributorsSignificance=0.07;ITSVertexerParam.suppressLowMultDebris=13;ITSVertexerParam.fineZWindow=0.013014502308220061;ITSVertexerParam.fineMinDensity=4;ITSVertexerParam.fineMaxDrift=0.004003953018135234;ITSVertexerParam.duplicateZScale=0.1230148343342053;ITSVertexerParam.duplicateZCut=0.06571512005272727;