Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

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

Merged
merged 4 commits into from Aug 2, 2022
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
a9cb462
Clamp the number of tracks to the capacity of the SoA
fwyzard Jul 28, 2022
7faa2a2
Check that the fixed-size SoA does not overflow
fwyzard Jul 28, 2022
9ebe6f3
Warn if the number of pixel track candidates was truncated
fwyzard Jul 30, 2022
d04920f
Simplify the final emplace
fwyzard Jul 30, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
11 changes: 9 additions & 2 deletions RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc
View file
Open in desktop
Expand Up @@ -60,10 +60,17 @@ void PixelTrackSoAFromCUDA::acquire(edm::Event const& iEvent,
}

void PixelTrackSoAFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& iSetup) {
#ifdef PIXEL_DEBUG_PRODUCE
// check that the fixed-size SoA does not overflow
auto const& tsoa = *soa_;
auto maxTracks = tsoa.stride();
auto nTracks = tsoa.nTracks();
assert(nTracks < maxTracks);
if (nTracks == maxTracks - 1) {
edm::LogWarning("PixelTracks") << "Unsorted reconstructed pixel tracks truncated to " << maxTracks - 1
<< " candidates";
}

#ifdef PIXEL_DEBUG_PRODUCE
std::cout << "size of SoA " << sizeof(tsoa) << " stride " << maxTracks << std::endl;
std::cout << "found " << nTracks << " tracks in cpu SoA at " << &tsoa << std::endl;

Expand All @@ -79,7 +86,7 @@ void PixelTrackSoAFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& i
#endif

// DO NOT make a copy (actually TWO....)
iEvent.emplace(tokenSOA_, PixelTrackHeterogeneous(std::move(soa_)));
iEvent.emplace(tokenSOA_, std::move(soa_));

assert(!soa_);
}
Expand Down
3 changes: 2 additions & 1 deletion RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h
View file
Open in desktop
Expand Up @@ -562,7 +562,8 @@ __global__ void kernel_fillHitDetIndices(HitContainer const *__restrict__ tuples
__global__ void kernel_fillNLayers(TkSoA *__restrict__ ptracks, cms::cuda::AtomicPairCounter *apc) {
auto &tracks = *ptracks;
auto first = blockIdx.x * blockDim.x + threadIdx.x;
auto ntracks = apc->get().m;
// clamp the number of tracks to the capacity of the SoA
auto ntracks = std::min<int>(apc->get().m, tracks.stride() - 1);
if (0 == first)
tracks.setNTracks(ntracks);
for (int idx = first, nt = ntracks; idx < nt; idx += gridDim.x * blockDim.x) {
Expand Down
10 changes: 10 additions & 0 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc
View file
Open in desktop
Expand Up @@ -247,5 +247,15 @@ PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuples(TrackingRecHit2DC
std::cout << "finished building pixel tracks on CPU" << std::endl;
#endif

// check that the fixed-size SoA does not overflow
auto const& tsoa = *soa;
auto maxTracks = tsoa.stride();
auto nTracks = tsoa.nTracks();
assert(nTracks < maxTracks);
if (nTracks == maxTracks - 1) {
edm::LogWarning("PixelTracks") << "Unsorted reconstructed pixel tracks truncated to " << maxTracks - 1
<< " candidates";
}

return tracks;
}