Clamp the number of pixel tracks to the capacity of the SoA #38883

fwyzard · 2022-07-28T15:37:38Z

PR description:

The data structure used to store the pixel tracks currently has a fixed size of 32768, which was found to be adequate based on MC studies.
While running over collisions data a small fraction of events has been found to have a larger number of tracks, leading to assertion failures or memory access errors.

This PR clamps the number of pixel tracks to the capacity of the SoA, and adds a warning like

%MSG-w PixelTracks:  PixelTrackSoAFromCUDA:hltPixelTracksFromGPU  30-Jul-2022 09:16:12 CEST Run: 356071 Event: 46424286
Unsorted reconstructed pixel tracks truncated to 32767 candidates
%MSG

PR validation:

Running over an event with > 33k pixel tracks no longer causes assertion failures.

If this PR will be backported please specify to which release cycle the backport is meant for:

To be backported to CMSSW 12.4.x for data taking.

fwyzard · 2022-07-28T15:38:15Z

type bugfix

fwyzard · 2022-07-28T15:38:30Z

enable gpu

fwyzard · 2022-07-28T15:38:35Z

please test

fwyzard · 2022-07-28T15:42:09Z

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h

+  uint32_t first = blockIdx.x * blockDim.x + threadIdx.x;
+  // clamp the number of tracks to the capacity of the SoA
+  auto ntracks = std::min<uint32_t>(apc->get().m, tracks.stride() - 1);


@VinInn does this look like a good fix to you ?
An alternative could be to adapt all client code to check the size vs the stride (or look for the nhits == 0 marker), but this seemed easier to implement and safer to use.

On the other this will completely silence the error.
Maybe we could check if nTracks() == stride() - 1 and assume it to mean that an "overflow" occurred ?

cmsbuild · 2022-07-28T15:45:29Z

+code-checks

Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-38883/31314

This PR adds an extra 24KB to repository
There are other open Pull requests which might conflict with changes you have proposed:
- File RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc modified in PR(s): Tracker Traits and Enabling Phase2 for Inner Tracker Reconstruction on GPU #38761, New Heterogeneous Memory Pool #37952
- File RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h modified in PR(s): Tracker Traits and Enabling Phase2 for Inner Tracker Reconstruction on GPU #38761, Use cooperative groups to populate Associations (Histograms) in Pixel Patatrack #35713

cmsbuild · 2022-07-28T15:45:53Z

A new Pull Request was created by @fwyzard (Andrea Bocci) for master.

It involves the following packages:

RecoPixelVertexing/PixelTrackFitting (reconstruction)
RecoPixelVertexing/PixelTriplets (reconstruction)

@jpata, @clacaputo can you please review it and eventually sign? Thanks.
@VourMa, @felicepantaleo, @GiacomoSguazzoni, @JanFSchulte, @rovere, @VinInn, @mmusich, @mtosi, @dgulhan this is something you requested to watch as well.
@perrotta, @dpiparo, @qliphy, @rappoccio you are the release manager for this.

cms-bot commands are listed here

VinInn · 2022-07-28T16:37:18Z

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h

  if (0 == first)
    tracks.setNTracks(ntracks);
-  for (int idx = first, nt = ntracks; idx < nt; idx += gridDim.x * blockDim.x) {
+  for (uint32_t idx = first; idx < ntracks; idx += gridDim.x * blockDim.x) {
    auto nHits = tracks.nHits(idx);


why?
int was intentional.

which change ?

nt = ntracks: I don't see how it can make a difference (ntracks is also a local variable) but I can revert it

int vs uint32_t: becuase ntracks is now uint32_t and the compiler warns about signed-vs-unsigned comparison

cmsbuild · 2022-07-28T19:57:05Z

+1

Summary: https://cmssdt.cern.ch/SDT/jenkins-artifacts/pull-request-integration/PR-197dc3/26510/summary.html
COMMIT: 9c1a5dc
CMSSW: CMSSW_12_5_X_2022-07-28-1100/el8_amd64_gcc10
Additional Tests: GPU
User test area: For local testing, you can use /cvmfs/cms-ci.cern.ch/week1/cms-sw/cmssw/38883/26510/install.sh to create a dev area with all the needed externals and cmssw changes.

GPU Comparison Summary

Summary:

No significant changes to the logs found
Reco comparison results: 0 differences found in the comparisons
DQMHistoTests: Total files compared: 4
DQMHistoTests: Total histograms compared: 19876
DQMHistoTests: Total failures: 36
DQMHistoTests: Total nulls: 0
DQMHistoTests: Total successes: 19840
DQMHistoTests: Total skipped: 0
DQMHistoTests: Total Missing objects: 0
DQMHistoSizes: Histogram memory added: 0.0 KiB( 3 files compared)
Checked 12 log files, 9 edm output root files, 4 DQM output files
TriggerResults: no differences found

Comparison Summary

Summary:

No significant changes to the logs found
Reco comparison results: 6 differences found in the comparisons
DQMHistoTests: Total files compared: 51
DQMHistoTests: Total histograms compared: 3668050
DQMHistoTests: Total failures: 8
DQMHistoTests: Total nulls: 0
DQMHistoTests: Total successes: 3668020
DQMHistoTests: Total skipped: 22
DQMHistoTests: Total Missing objects: 0
DQMHistoSizes: Histogram memory added: 0.0 KiB( 50 files compared)
Checked 210 log files, 47 edm output root files, 51 DQM output files
TriggerResults: no differences found

clacaputo · 2022-07-29T10:21:17Z

+reconstruction

cmsbuild · 2022-07-29T10:21:40Z

This pull request is fully signed and it will be integrated in one of the next master IBs (tests are also fine). This pull request will now be reviewed by the release team before it's merged. @perrotta, @dpiparo, @qliphy, @rappoccio (and backports should be raised in the release meeting by the corresponding L2)

fwyzard · 2022-07-29T11:50:42Z

hold

I'm waiting for further feedback from @VinInn and mulling on the possibility of adding a LogWarning.

cmsbuild · 2022-07-29T11:51:06Z

Pull request has been put on hold by @fwyzard
They need to issue an unhold command to remove the hold state or L1 can unhold it for all

fwyzard · 2022-07-30T08:00:20Z

please test

cmsbuild · 2022-07-30T08:07:28Z

+code-checks

Logs: https://cmssdt.cern.ch/SDT/code-checks/cms-sw-PR-38883/31340

This PR adds an extra 28KB to repository
There are other open Pull requests which might conflict with changes you have proposed:
- File RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc modified in PR(s): Tracker Traits and Enabling Phase2 for Inner Tracker Reconstruction on GPU #38761, New Heterogeneous Memory Pool #37952
- File RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h modified in PR(s): Tracker Traits and Enabling Phase2 for Inner Tracker Reconstruction on GPU #38761, Use cooperative groups to populate Associations (Histograms) in Pixel Patatrack #35713
- File RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc modified in PR(s): Tracker Traits and Enabling Phase2 for Inner Tracker Reconstruction on GPU #38761, New Heterogeneous Memory Pool #37952

cmsbuild · 2022-07-30T08:07:51Z

Pull request #38883 was updated. @jpata, @clacaputo can you please check and sign again.

cmsbuild · 2022-07-30T12:11:34Z

+1

Summary: https://cmssdt.cern.ch/SDT/jenkins-artifacts/pull-request-integration/PR-197dc3/26543/summary.html
COMMIT: d04920f
CMSSW: CMSSW_12_5_X_2022-07-29-2300/el8_amd64_gcc10
Additional Tests: GPU
User test area: For local testing, you can use /cvmfs/cms-ci.cern.ch/week1/cms-sw/cmssw/38883/26543/install.sh to create a dev area with all the needed externals and cmssw changes.

GPU Comparison Summary

Summary:

No significant changes to the logs found
Reco comparison results: 0 differences found in the comparisons
DQMHistoTests: Total files compared: 4
DQMHistoTests: Total histograms compared: 19876
DQMHistoTests: Total failures: 36
DQMHistoTests: Total nulls: 0
DQMHistoTests: Total successes: 19840
DQMHistoTests: Total skipped: 0
DQMHistoTests: Total Missing objects: 0
DQMHistoSizes: Histogram memory added: 0.0 KiB( 3 files compared)
Checked 12 log files, 9 edm output root files, 4 DQM output files
TriggerResults: found differences in 1 / 3 workflows

Comparison Summary

Summary:

No significant changes to the logs found
Reco comparison results: 8 differences found in the comparisons
DQMHistoTests: Total files compared: 51
DQMHistoTests: Total histograms compared: 3668050
DQMHistoTests: Total failures: 13
DQMHistoTests: Total nulls: 1
DQMHistoTests: Total successes: 3668014
DQMHistoTests: Total skipped: 22
DQMHistoTests: Total Missing objects: 0
DQMHistoSizes: Histogram memory added: 0.004 KiB( 50 files compared)
DQMHistoSizes: changed ( 312.0 ): 0.004 KiB MessageLogger/Warnings
Checked 210 log files, 47 edm output root files, 51 DQM output files
TriggerResults: no differences found

clacaputo · 2022-07-31T06:52:46Z

+reconstruction

cmsbuild · 2022-07-31T06:53:11Z

This pull request is fully signed and it will be integrated in one of the next master IBs (tests are also fine). This pull request will now be reviewed by the release team before it's merged. @perrotta, @dpiparo, @qliphy, @rappoccio (and backports should be raised in the release meeting by the corresponding L2)

fwyzard · 2022-08-01T07:57:37Z

urgent

qliphy · 2022-08-01T08:30:41Z

+1

cmsbuild added this to the CMSSW_12_5_X milestone Jul 28, 2022

cmsbuild added code-checks-pending orp-pending pending-signatures reconstruction-pending tests-pending labels Jul 28, 2022

fwyzard mentioned this pull request Jul 28, 2022

Clamp the number of pixel tracks to the capacity of the SoA [12.4.x] #38884

Merged

cmsbuild added the bug-fix label Jul 28, 2022

cmsbuild added tests-started and removed tests-pending labels Jul 28, 2022

fwyzard commented Jul 28, 2022

View reviewed changes

cmsbuild added code-checks-approved and removed code-checks-pending labels Jul 28, 2022

VinInn reviewed Jul 28, 2022

View reviewed changes

cmsbuild added tests-approved and removed tests-started labels Jul 28, 2022

cmsbuild added fully-signed reconstruction-approved and removed reconstruction-pending pending-signatures labels Jul 29, 2022

fwyzard added 2 commits July 30, 2022 09:59

Warn if the number of pixel track candidates was truncated

9ebe6f3

Simplify the final emplace

d04920f

fwyzard force-pushed the Clamp_the_number_of_pixel_tracks_to_the_capacity_of_the_SoA branch from f3c7494 to d04920f Compare July 30, 2022 07:59

cmsbuild added code-checks-pending tests-pending and removed tests-started code-checks-rejected labels Jul 30, 2022

cmsbuild added tests-started and removed tests-pending labels Jul 30, 2022

cmsbuild added code-checks-approved and removed code-checks-pending labels Jul 30, 2022

cmsbuild added tests-approved and removed tests-started labels Jul 30, 2022

cmsbuild added fully-signed reconstruction-approved and removed reconstruction-pending pending-signatures labels Jul 31, 2022

cmsbuild added the urgent label Aug 1, 2022

cmsbuild added orp-approved and removed orp-pending labels Aug 1, 2022

cmsbuild merged commit ce3448c into cms-sw:master Aug 1, 2022

fwyzard deleted the Clamp_the_number_of_pixel_tracks_to_the_capacity_of_the_SoA branch October 29, 2022 08:37

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Clamp the number of pixel tracks to the capacity of the SoA #38883

Clamp the number of pixel tracks to the capacity of the SoA #38883

fwyzard commented Jul 28, 2022 •

edited

Loading

fwyzard commented Jul 28, 2022

fwyzard commented Jul 28, 2022

fwyzard commented Jul 28, 2022

fwyzard Jul 28, 2022

cmsbuild commented Jul 28, 2022

cmsbuild commented Jul 28, 2022

VinInn Jul 28, 2022

fwyzard Jul 28, 2022

cmsbuild commented Jul 28, 2022

clacaputo commented Jul 29, 2022

cmsbuild commented Jul 29, 2022

fwyzard commented Jul 29, 2022

cmsbuild commented Jul 29, 2022

fwyzard commented Jul 30, 2022

cmsbuild commented Jul 30, 2022

cmsbuild commented Jul 30, 2022

cmsbuild commented Jul 30, 2022

clacaputo commented Jul 31, 2022

cmsbuild commented Jul 31, 2022

fwyzard commented Aug 1, 2022

qliphy commented Aug 1, 2022

Clamp the number of pixel tracks to the capacity of the SoA #38883

Clamp the number of pixel tracks to the capacity of the SoA #38883

Conversation

fwyzard commented Jul 28, 2022 • edited Loading

PR description:

PR validation:

If this PR will be backported please specify to which release cycle the backport is meant for:

fwyzard commented Jul 28, 2022

fwyzard commented Jul 28, 2022

fwyzard commented Jul 28, 2022

fwyzard Jul 28, 2022

Choose a reason for hiding this comment

cmsbuild commented Jul 28, 2022

cmsbuild commented Jul 28, 2022

VinInn Jul 28, 2022

Choose a reason for hiding this comment

fwyzard Jul 28, 2022

Choose a reason for hiding this comment

cmsbuild commented Jul 28, 2022

GPU Comparison Summary

Comparison Summary

clacaputo commented Jul 29, 2022

cmsbuild commented Jul 29, 2022

fwyzard commented Jul 29, 2022

cmsbuild commented Jul 29, 2022

fwyzard commented Jul 30, 2022

cmsbuild commented Jul 30, 2022

cmsbuild commented Jul 30, 2022

cmsbuild commented Jul 30, 2022

GPU Comparison Summary

Comparison Summary

clacaputo commented Jul 31, 2022

cmsbuild commented Jul 31, 2022

fwyzard commented Aug 1, 2022

qliphy commented Aug 1, 2022

fwyzard commented Jul 28, 2022 •

edited

Loading