From 07c25b655816834b8242575bcb6211927496a45d Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 6 Apr 2021 15:09:00 +0200 Subject: [PATCH 01/10] Fixed include guards naming. (#614) https://github.com/cms-sw/cmssw/pull/31722#pullrequestreview-623346747 --- CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h index f74717c41e4d7..a827214cf8222 100644 --- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h +++ b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h @@ -1,5 +1,5 @@ -#ifndef CUDADataFormats_Track_TrackHeterogeneousT_H -#define CUDADataFormats_Track_TrackHeterogeneousT_H +#ifndef CUDADataFormats_Track_PixelTrackHeterogeneousT_h +#define CUDADataFormats_Track_PixelTrackHeterogeneousT_h #include "CUDADataFormats/Track/interface/TrajectoryStateSoAT.h" #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h" @@ -70,4 +70,4 @@ namespace pixelTrack { } // namespace pixelTrack -#endif // CUDADataFormats_Track_TrackHeterogeneousT_H +#endif // CUDADataFormats_Track_PixelTrackHeterogeneousT_h From 7467e172968342a6091bff79cd7685763d1718b1 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 6 Apr 2021 15:15:22 +0200 Subject: [PATCH 02/10] Unified member access methods for class TrackSoAHeterogeneousT. (#614) Used accessors in all cases. This commit does not change the memory layout of the class. Also fixed C-style casts (reinterpret casts are needed here). https://github.com/cms-sw/cmssw/pull/31722#discussion_r603382606 https://github.com/cms-sw/cmssw/pull/31722#discussion_r603383530 --- .../Track/interface/TrackSoAHeterogeneousT.h | 70 ++++++++++++------- .../plugins/PixelTrackProducerFromSoA.cc | 4 +- .../plugins/BrokenLineFitOnGPU.h | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cc | 6 +- .../plugins/CAHitNtupletGeneratorKernels.cu | 6 +- .../CAHitNtupletGeneratorKernelsImpl.h | 2 +- .../plugins/CAHitNtupletGeneratorOnGPU.cc | 4 +- .../PixelTriplets/plugins/RiemannFitOnGPU.h | 2 +- .../plugins/SeedProducerFromSoA.cc | 4 +- 9 files changed, 61 insertions(+), 39 deletions(-) diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h index a827214cf8222..8daf780644749 100644 --- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h +++ b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h @@ -19,39 +19,61 @@ class TrackSoAHeterogeneousT { using hindex_type = uint32_t; using HitContainer = cms::cuda::OneToManyAssoc; + + // quality accessors + constexpr Quality quality(int32_t i) const { return reinterpret_cast(quality_(i)); } + constexpr Quality &quality(int32_t i) { return reinterpret_cast(quality_(i)); } + constexpr Quality const *qualityData() const { return reinterpret_cast(quality_.data()); } + constexpr Quality *qualityData() { return reinterpret_cast(quality_.data()); } + + // chi2 accessors + constexpr auto & chi2(int32_t i) { return chi2_(i); } + constexpr auto chi2(int32_t i) const { return chi2_(i); } + + // stateAtBS accessors + constexpr auto & stateAtBS() { return stateAtBS_; } + constexpr auto stateAtBS() const { return stateAtBS_; } + // eta accessors + constexpr auto & eta(int32_t i) { return eta_(i); } + constexpr auto eta(int32_t i) const { return eta_(i); } + // pt accessors + constexpr auto & pt(int32_t i) { return pt_(i); } + constexpr auto pt(int32_t i) const { return pt_(i); } + + constexpr float charge(int32_t i) const { return std::copysign(1.f, stateAtBS_.state(i)(2)); } + constexpr float phi(int32_t i) const { return stateAtBS_.state(i)(0); } + constexpr float tip(int32_t i) const { return stateAtBS_.state(i)(1); } + constexpr float zip(int32_t i) const { return stateAtBS_.state(i)(4); } + + // hitIndices accessors + constexpr auto & hitIndices() { return hitIndices_; } + constexpr auto const & hitIndices() const { return hitIndices_; } + + // detInndices accessor + constexpr int nHits(int i) const { return detIndices_.size(i); } + constexpr auto & detIndices() { return detIndices_; } + constexpr auto const & detIndices() const { return detIndices_; } + + // state at the detector of the outermost hit + // representation to be decided... + // not yet filled on GPU + // TrajectoryStateSoA stateAtOuterDet; +private: // Always check quality is at least loose! // CUDA does not support enums in __lgc ... -private: eigenSoA::ScalarSoA quality_; -public: - constexpr Quality quality(int32_t i) const { return (Quality)(quality_(i)); } - constexpr Quality &quality(int32_t i) { return (Quality &)(quality_(i)); } - constexpr Quality const *qualityData() const { return (Quality const *)(quality_.data()); } - constexpr Quality *qualityData() { return (Quality *)(quality_.data()); } - // this is chi2/ndof as not necessarely all hits are used in the fit - eigenSoA::ScalarSoA chi2; - - constexpr int nHits(int i) const { return detIndices.size(i); } + eigenSoA::ScalarSoA chi2_; // State at the Beam spot // phi,tip,1/pt,cotan(theta),zip - TrajectoryStateSoAT stateAtBS; - eigenSoA::ScalarSoA eta; - eigenSoA::ScalarSoA pt; - constexpr float charge(int32_t i) const { return std::copysign(1.f, stateAtBS.state(i)(2)); } - constexpr float phi(int32_t i) const { return stateAtBS.state(i)(0); } - constexpr float tip(int32_t i) const { return stateAtBS.state(i)(1); } - constexpr float zip(int32_t i) const { return stateAtBS.state(i)(4); } - - // state at the detector of the outermost hit - // representation to be decided... - // not yet filled on GPU - // TrajectoryStateSoA stateAtOuterDet; + TrajectoryStateSoAT stateAtBS_; + eigenSoA::ScalarSoA eta_; + eigenSoA::ScalarSoA pt_; - HitContainer hitIndices; - HitContainer detIndices; + HitContainer hitIndices_; + HitContainer detIndices_; }; namespace pixelTrack { diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index 60225eceebc00..ab57dfd1a0b69 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc @@ -133,8 +133,8 @@ void PixelTrackProducerFromSoA::produce(edm::StreamID streamID, const auto &tsoa = *iEvent.get(tokenTrack_); auto const *quality = tsoa.qualityData(); - auto const &fit = tsoa.stateAtBS; - auto const &hitIndices = tsoa.hitIndices; + auto const &fit = tsoa.stateAtBS(); + auto const &hitIndices = tsoa.hitIndices(); auto maxTracks = tsoa.stride(); int32_t nt = 0; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h index 87e93e6f0bd05..c932fe7f7fcf0 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h @@ -155,7 +155,7 @@ __global__ void kernel_BLFit(caConstants::TupleMultiplicity const *__restrict__ brokenline::lineFit(hits_ge, fast_fit, bField, data, line); brokenline::circleFit(hits, hits_ge, fast_fit, bField, data, circle); - results->stateAtBS.copyFromCircle(circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid); + results->stateAtBS().copyFromCircle(circle.par, circle.cov, line.par, line.cov, 1.f / float(bField), tkid); results->pt(tkid) = float(bField) / float(std::abs(circle.par(2))); results->eta(tkid) = asinhf(line.par(0)); results->chi2(tkid) = (circle.chi2 + line.chi2) / (2 * N - 5); diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc index 5a2157bf87b7c..cf54aab22508d 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc @@ -7,7 +7,7 @@ void CAHitNtupletGeneratorKernelsCPU::printCounters(Counters const *counters) { template <> void CAHitNtupletGeneratorKernelsCPU::fillHitDetIndices(HitsView const *hv, TkSoA *tracks_d, cudaStream_t) { - kernel_fillHitDetIndices(&tracks_d->hitIndices, hv, &tracks_d->detIndices); + kernel_fillHitDetIndices(&tracks_d->hitIndices(), hv, &tracks_d->detIndices()); } template <> @@ -75,7 +75,7 @@ void CAHitNtupletGeneratorKernelsCPU::buildDoublets(HitsOnCPU const &hh, cudaStr template <> void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) { - auto *tuples_d = &tracks_d->hitIndices; + auto *tuples_d = &tracks_d->hitIndices(); auto *quality_d = tracks_d->qualityData(); assert(tuples_d && quality_d); @@ -154,7 +154,7 @@ void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA * template <> void CAHitNtupletGeneratorKernelsCPU::classifyTuples(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) { - auto const *tuples_d = &tracks_d->hitIndices; + auto const *tuples_d = &tracks_d->hitIndices(); auto *quality_d = tracks_d->qualityData(); // classify tracks based on kinematics diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu index ca334223f8fc7..2da78a2adec81 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu @@ -6,7 +6,7 @@ void CAHitNtupletGeneratorKernelsGPU::fillHitDetIndices(HitsView const *hv, TkSo auto numberOfBlocks = (HitContainer::ctCapacity() + blockSize - 1) / blockSize; kernel_fillHitDetIndices<<>>( - &tracks_d->hitIndices, hv, &tracks_d->detIndices); + &tracks_d->hitIndices(), hv, &tracks_d->detIndices()); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG cudaDeviceSynchronize(); @@ -17,7 +17,7 @@ void CAHitNtupletGeneratorKernelsGPU::fillHitDetIndices(HitsView const *hv, TkSo template <> void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) { // these are pointer on GPU! - auto *tuples_d = &tracks_d->hitIndices; + auto *tuples_d = &tracks_d->hitIndices(); auto *quality_d = tracks_d->qualityData(); // zero tuples @@ -224,7 +224,7 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(HitsOnCPU const &hh, cudaStr template <> void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA *tracks_d, cudaStream_t cudaStream) { // these are pointer on GPU! - auto const *tuples_d = &tracks_d->hitIndices; + auto const *tuples_d = &tracks_d->hitIndices(); auto *quality_d = tracks_d->qualityData(); int32_t nhits = hh.nHits(); diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h index 550541fdca6fb..a54d75716c2d2 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h @@ -361,7 +361,7 @@ __global__ void kernel_classifyTracks(HitContainer const *__restrict__ tuples, // if the fit has any invalid parameters, mark it as bad bool isNaN = false; for (int i = 0; i < 5; ++i) { - isNaN |= std::isnan(tracks->stateAtBS.state(it)(i)); + isNaN |= std::isnan(tracks->stateAtBS().state(it)(i)); } if (isNaN) { #ifdef NTUPLE_DEBUG diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc index 3f5ba1d04e7db..f8ee734f3cf45 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc @@ -192,7 +192,7 @@ PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuplesAsync(TrackingRecH kernels.fillHitDetIndices(hits_d.view(), soa, stream); // in principle needed only if Hits not "available" HelixFitOnGPU fitter(bfield, m_params.fit5as4_); - fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa); + fitter.allocateOnGPU(&(soa->hitIndices()), kernels.tupleMultiplicity(), soa); if (m_params.useRiemannFit_) { fitter.launchRiemannKernels(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets, stream); } else { @@ -228,7 +228,7 @@ PixelTrackHeterogeneous CAHitNtupletGeneratorOnGPU::makeTuples(TrackingRecHit2DC // now fit HelixFitOnGPU fitter(bfield, m_params.fit5as4_); - fitter.allocateOnGPU(&(soa->hitIndices), kernels.tupleMultiplicity(), soa); + fitter.allocateOnGPU(&(soa->hitIndices()), kernels.tupleMultiplicity(), soa); if (m_params.useRiemannFit_) { fitter.launchRiemannKernelsOnCPU(hits_d.view(), hits_d.nHits(), caConstants::maxNumberOfQuadruplets); diff --git a/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h index 926002d674b83..2ebb544203156 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h @@ -159,7 +159,7 @@ __global__ void kernel_LineFit(caConstants::TupleMultiplicity const *__restrict_ riemannFit::fromCircleToPerigee(circle_fit[local_idx]); - results->stateAtBS.copyFromCircle( + results->stateAtBS().copyFromCircle( circle_fit[local_idx].par, circle_fit[local_idx].cov, line_fit.par, line_fit.cov, 1.f / float(bField), tkid); results->pt(tkid) = bField / std::abs(circle_fit[local_idx].par(2)); results->eta(tkid) = asinhf(line_fit.par(0)); diff --git a/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc b/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc index 0e5823fc46c46..fd06cb4d80ac6 100644 --- a/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc +++ b/RecoTracker/TkSeedGenerator/plugins/SeedProducerFromSoA.cc @@ -92,8 +92,8 @@ void SeedProducerFromSoA::produce(edm::StreamID streamID, edm::Event& iEvent, co const auto& tsoa = *(iEvent.get(tokenTrack_)); auto const* quality = tsoa.qualityData(); - auto const& fit = tsoa.stateAtBS; - auto const& detIndices = tsoa.detIndices; + auto const& fit = tsoa.stateAtBS(); + auto const& detIndices = tsoa.detIndices(); auto maxTracks = tsoa.stride(); int32_t nt = 0; From 7222e8a7b806e12021fbf70bdbf2acb949d0dc84 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 6 Apr 2021 15:48:04 +0200 Subject: [PATCH 03/10] Fixed indlude guards naming. (#614) https://github.com/cms-sw/cmssw/pull/31722#discussion_r603384652 --- CUDADataFormats/Track/interface/TrajectoryStateSoAT.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h b/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h index 64fcd573a6991..958741e024c68 100644 --- a/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h +++ b/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h @@ -1,5 +1,5 @@ -#ifndef CUDADataFormats_Track_TrajectoryStateSOAT_H -#define CUDADataFormats_Track_TrajectoryStateSOAT_H +#ifndef CUDADataFormats_Track_TrajectoryStateSoAT_h +#define CUDADataFormats_Track_TrajectoryStateSoAT_h #include #include "HeterogeneousCore/CUDAUtilities/interface/eigenSoA.h" @@ -56,4 +56,4 @@ struct TrajectoryStateSoAT { } }; -#endif // CUDADataFormats_Track_TrajectoryStateSOAT_H +#endif // CUDADataFormats_Track_TrajectoryStateSoAT_h From c52e7323d9b1e8133313fc47cb7f8f876b9b9679 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 6 Apr 2021 16:21:19 +0200 Subject: [PATCH 04/10] Fixed C-style cast. (#614) https://github.com/cms-sw/cmssw/pull/31722#discussion_r603386189 --- CUDADataFormats/Track/interface/TrajectoryStateSoAT.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h b/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h index 958741e024c68..947b2af7aabb2 100644 --- a/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h +++ b/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h @@ -25,12 +25,12 @@ struct TrajectoryStateSoAT { auto cov = covariance(i); cov(0) = ccov(0, 0); cov(1) = ccov(0, 1); - cov(2) = b * float(ccov(0, 2)); + cov(2) = b * static_cast(ccov(0, 2)); cov(4) = cov(3) = 0; cov(5) = ccov(1, 1); - cov(6) = b * float(ccov(1, 2)); + cov(6) = b * static_cast(ccov(1, 2)); cov(8) = cov(7) = 0; - cov(9) = b * b * float(ccov(2, 2)); + cov(9) = b * b * static_cast(ccov(2, 2)); cov(11) = cov(10) = 0; cov(12) = lcov(0, 0); cov(13) = lcov(0, 1); From 01043c8d6f2a66979c0fd9fb9edf94aa9911032e Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Thu, 8 Apr 2021 15:26:34 +0200 Subject: [PATCH 05/10] Updated comments after methods renaming. (#614) https://github.com/cms-sw/cmssw/pull/31722#discussion_r603400548. --- .../PixelTrackFitting/interface/BrokenLine.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/RecoPixelVertexing/PixelTrackFitting/interface/BrokenLine.h b/RecoPixelVertexing/PixelTrackFitting/interface/BrokenLine.h index 86fe6a278777c..65ce5d4c3dbf7 100644 --- a/RecoPixelVertexing/PixelTrackFitting/interface/BrokenLine.h +++ b/RecoPixelVertexing/PixelTrackFitting/interface/BrokenLine.h @@ -535,9 +535,9 @@ namespace brokenline { /*! \brief Helix fit by three step: - -fast pre-fit (see Fast_fit() for further info); \n - -circle fit of the hits projected in the transverse plane by Broken Line algorithm (see BL_Circle_fit() for further info); \n - -line fit of the hits projected on the (pre-fitted) cilinder surface by Broken Line algorithm (see BL_Line_fit() for further info); \n + -fast pre-fit (see fastFit() for further info); \n + -circle fit of the hits projected in the transverse plane by Broken Line algorithm (see circleFit() for further info); \n + -line fit of the hits projected on the (pre-fitted) cilinder surface by Broken Line algorithm (see lineFit() for further info); \n Points must be passed ordered (from inner to outer layer). \param hits Matrix3xNd hits coordinates in this form: \n @@ -561,9 +561,9 @@ namespace brokenline { |(x1,z4)|(x2,z4)|(x3,z4)|(x4,z4)|.|(y1,z4)|(y2,z4)|(y3,z4)|(y4,z4)|.|(z1,z4)|(z2,z4)|(z3,z4)|(z4,z4)| \param bField magnetic field in the center of the detector in Gev/cm/c, in order to perform the p_t calculation. - \warning see BL_Circle_fit(), BL_Line_fit() and Fast_fit() warnings. + \warning see circleFit(), lineFit() and fastFit() warnings. - \bug see BL_Circle_fit(), BL_Line_fit() and Fast_fit() bugs. + \bug see circleFit(), lineFit() and fastFit() bugs. \return (phi,Tip,p_t,cot(theta)),Zip), their covariance matrix and the chi2's of the circle and line fits. */ From e783810b428edd8f2e64b1946afaf5ecd25f9c58 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Thu, 8 Apr 2021 16:28:26 +0200 Subject: [PATCH 06/10] Modified PixelTrackProducer class to use esConsume(). (#614) https://github.com/cms-sw/cmssw/pull/31722#discussion_r603404617 --- .../PixelTrackFitting/plugins/PixelTrackProducer.cc | 9 +++++---- .../PixelTrackFitting/plugins/PixelTrackProducer.h | 2 ++ 2 files changed, 7 insertions(+), 4 deletions(-) diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc index 91c3a44cc8643..a276af79b0945 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc @@ -21,7 +21,9 @@ using namespace pixeltrackfitting; using edm::ParameterSet; -PixelTrackProducer::PixelTrackProducer(const ParameterSet& cfg) : theReconstruction(cfg, consumesCollector()) { +PixelTrackProducer::PixelTrackProducer(const ParameterSet& cfg) + : theReconstruction(cfg, consumesCollector()), + htTopoToken_(esConsumes()) { edm::LogInfo("PixelTrackProducer") << " construction..."; produces(); produces(); @@ -44,9 +46,8 @@ void PixelTrackProducer::produce(edm::Event& ev, const edm::EventSetup& es) { TracksWithTTRHs tracks; theReconstruction.run(tracks, ev, es); - edm::ESHandle httopo; - es.get().get(httopo); + auto htTopo = es.getData(htTopoToken_); // store tracks - storeTracks(ev, tracks, *httopo); + storeTracks(ev, tracks, htTopo); } diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h index c38fd44c0d7f5..0d492f979f3aa 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h @@ -3,6 +3,7 @@ #include "FWCore/Framework/interface/stream/EDProducer.h" #include "RecoPixelVertexing/PixelTrackFitting/interface/PixelTrackReconstruction.h" +#include "Geometry/Records/interface/TrackerTopologyRcd.h" namespace edm { class Event; @@ -24,6 +25,7 @@ class PixelTrackProducer : public edm::stream::EDProducer<> { private: PixelTrackReconstruction theReconstruction; + const edm::ESGetToken htTopoToken_; }; #endif // RecoPixelVertexing_PixelTrackFitting_plugins_PixelTrackProducer_h From 323eca77ac7e365f96eca5392587434598c9cb54 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Mon, 12 Apr 2021 11:39:30 +0200 Subject: [PATCH 07/10] Returned complex strcuture by reference instead of value. Also made the type explicit. https://github.com/cms-patatrack/cmssw/pull/615#discussion_r611021950 --- CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h index 8daf780644749..003d1d184f845 100644 --- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h +++ b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h @@ -31,8 +31,8 @@ class TrackSoAHeterogeneousT { constexpr auto chi2(int32_t i) const { return chi2_(i); } // stateAtBS accessors - constexpr auto & stateAtBS() { return stateAtBS_; } - constexpr auto stateAtBS() const { return stateAtBS_; } + constexpr TrajectoryStateSoAT & stateAtBS() { return stateAtBS_; } + constexpr TrajectoryStateSoAT const & stateAtBS() const { return stateAtBS_; } // eta accessors constexpr auto & eta(int32_t i) { return eta_(i); } constexpr auto eta(int32_t i) const { return eta_(i); } From 4ed484d17b78aea6e64578cff4f3b2c73d82bd1b Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Tue, 13 Apr 2021 09:08:31 +0200 Subject: [PATCH 08/10] Fused class, header and SealModule.cc file. https://github.com/cms-sw/cmssw/pull/31722#discussion_r603659618 --- .../plugins/PixelTrackProducer.cc | 70 ++++++++++++------- .../plugins/PixelTrackProducer.h | 31 -------- .../PixelTrackFitting/plugins/SealModule.cc | 5 -- 3 files changed, 46 insertions(+), 60 deletions(-) delete mode 100644 RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h delete mode 100644 RecoPixelVertexing/PixelTrackFitting/plugins/SealModule.cc diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc index a276af79b0945..87220633c9a43 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc @@ -15,39 +15,61 @@ #include "FWCore/ParameterSet/interface/ParameterSetDescription.h" #include "Geometry/Records/interface/TrackerTopologyRcd.h" -#include "PixelTrackProducer.h" +#include "FWCore/Framework/interface/stream/EDProducer.h" +#include "RecoPixelVertexing/PixelTrackFitting/interface/PixelTrackReconstruction.h" +#include "Geometry/Records/interface/TrackerTopologyRcd.h" #include "storeTracks.h" +#include "FWCore/PluginManager/interface/ModuleDef.h" +#include "FWCore/Framework/interface/MakerMacros.h" + +namespace edm { + class Event; + class EventSetup; + class ParameterSet; + class ConfigurationDescriptions; +} // namespace edm +class TrackerTopology; + using namespace pixeltrackfitting; using edm::ParameterSet; -PixelTrackProducer::PixelTrackProducer(const ParameterSet& cfg) - : theReconstruction(cfg, consumesCollector()), - htTopoToken_(esConsumes()) { - edm::LogInfo("PixelTrackProducer") << " construction..."; - produces(); - produces(); - produces(); -} +class PixelTrackProducer : public edm::stream::EDProducer<> { +public: + explicit PixelTrackProducer(const edm::ParameterSet& cfg) + : theReconstruction(cfg, consumesCollector()), + htTopoToken_(esConsumes()) { + edm::LogInfo("PixelTrackProducer") << " construction..."; + produces(); + produces(); + produces(); + } + + ~PixelTrackProducer() override {} + + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; -PixelTrackProducer::~PixelTrackProducer() {} + desc.add("passLabel", "pixelTracks"); // What is this? It is not used anywhere in this code. + PixelTrackReconstruction::fillDescriptions(desc); -void PixelTrackProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { - edm::ParameterSetDescription desc; + descriptions.add("pixelTracks", desc); + } - desc.add("passLabel", "pixelTracks"); // What is this? It is not used anywhere in this code. - PixelTrackReconstruction::fillDescriptions(desc); + void produce(edm::Event& ev, const edm::EventSetup& es) override { + LogDebug("PixelTrackProducer, produce") << "event# :" << ev.id(); - descriptions.add("pixelTracks", desc); -} + TracksWithTTRHs tracks; + theReconstruction.run(tracks, ev, es); + auto htTopo = es.getData(htTopoToken_); -void PixelTrackProducer::produce(edm::Event& ev, const edm::EventSetup& es) { - LogDebug("PixelTrackProducer, produce") << "event# :" << ev.id(); + // store tracks + storeTracks(ev, tracks, htTopo); + } - TracksWithTTRHs tracks; - theReconstruction.run(tracks, ev, es); - auto htTopo = es.getData(htTopoToken_); +private: + PixelTrackReconstruction theReconstruction; + const edm::ESGetToken htTopoToken_; +}; - // store tracks - storeTracks(ev, tracks, htTopo); -} +DEFINE_FWK_MODULE(PixelTrackProducer); diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h deleted file mode 100644 index 0d492f979f3aa..0000000000000 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h +++ /dev/null @@ -1,31 +0,0 @@ -#ifndef RecoPixelVertexing_PixelTrackFitting_plugins_PixelTrackProducer_h -#define RecoPixelVertexing_PixelTrackFitting_plugins_PixelTrackProducer_h - -#include "FWCore/Framework/interface/stream/EDProducer.h" -#include "RecoPixelVertexing/PixelTrackFitting/interface/PixelTrackReconstruction.h" -#include "Geometry/Records/interface/TrackerTopologyRcd.h" - -namespace edm { - class Event; - class EventSetup; - class ParameterSet; - class ConfigurationDescriptions; -} // namespace edm -class TrackerTopology; - -class PixelTrackProducer : public edm::stream::EDProducer<> { -public: - explicit PixelTrackProducer(const edm::ParameterSet& conf); - - ~PixelTrackProducer() override; - - static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); - - void produce(edm::Event& ev, const edm::EventSetup& es) override; - -private: - PixelTrackReconstruction theReconstruction; - const edm::ESGetToken htTopoToken_; -}; - -#endif // RecoPixelVertexing_PixelTrackFitting_plugins_PixelTrackProducer_h diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/SealModule.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/SealModule.cc deleted file mode 100644 index 11a4c3e12c308..0000000000000 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/SealModule.cc +++ /dev/null @@ -1,5 +0,0 @@ -#include "FWCore/PluginManager/interface/ModuleDef.h" -#include "FWCore/Framework/interface/MakerMacros.h" - -#include "PixelTrackProducer.h" -DEFINE_FWK_MODULE(PixelTrackProducer); From 0eb287c0d01d3fddc236873a2f3a1bd91360ebcb Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Thu, 15 Apr 2021 11:23:09 +0200 Subject: [PATCH 09/10] Used a wider NaN test to also cover possible Inf with -ofast https://github.com/cms-sw/cmssw/pull/31722#discussion_r603678116 --- .../plugins/CAHitNtupletGeneratorKernelsImpl.h | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h index a54d75716c2d2..310208f442856 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h @@ -13,6 +13,7 @@ #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h" +#include "FWCore/Utilities/interface/isFinite.h" #include "CAConstants.h" #include "CAHitNtupletGeneratorKernels.h" @@ -359,11 +360,11 @@ __global__ void kernel_classifyTracks(HitContainer const *__restrict__ tuples, continue; // if the fit has any invalid parameters, mark it as bad - bool isNaN = false; + bool isNotFinite = false; for (int i = 0; i < 5; ++i) { - isNaN |= std::isnan(tracks->stateAtBS().state(it)(i)); + isNotFinite |= edm::isNotFinite(tracks->stateAtBS().state(it)(i)); } - if (isNaN) { + if (isNotFinite) { #ifdef NTUPLE_DEBUG printf("NaN in fit %d size %d chi2 %f\n", it, tuples->size(it), tracks->chi2(it)); #endif From e1541ce42fde5c4700a817c6e0f6b93dd55007f6 Mon Sep 17 00:00:00 2001 From: Eric Cano Date: Thu, 27 May 2021 15:47:20 +0200 Subject: [PATCH 10/10] Fixed code format --- .../Track/interface/TrackSoAHeterogeneousT.h | 19 +++++++++---------- .../plugins/PixelTrackProducer.cc | 3 +-- 2 files changed, 10 insertions(+), 12 deletions(-) diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h index 003d1d184f845..0f0f6d21e74fc 100644 --- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h +++ b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h @@ -19,7 +19,6 @@ class TrackSoAHeterogeneousT { using hindex_type = uint32_t; using HitContainer = cms::cuda::OneToManyAssoc; - // quality accessors constexpr Quality quality(int32_t i) const { return reinterpret_cast(quality_(i)); } constexpr Quality &quality(int32_t i) { return reinterpret_cast(quality_(i)); } @@ -27,17 +26,17 @@ class TrackSoAHeterogeneousT { constexpr Quality *qualityData() { return reinterpret_cast(quality_.data()); } // chi2 accessors - constexpr auto & chi2(int32_t i) { return chi2_(i); } + constexpr auto &chi2(int32_t i) { return chi2_(i); } constexpr auto chi2(int32_t i) const { return chi2_(i); } // stateAtBS accessors - constexpr TrajectoryStateSoAT & stateAtBS() { return stateAtBS_; } - constexpr TrajectoryStateSoAT const & stateAtBS() const { return stateAtBS_; } + constexpr TrajectoryStateSoAT &stateAtBS() { return stateAtBS_; } + constexpr TrajectoryStateSoAT const &stateAtBS() const { return stateAtBS_; } // eta accessors - constexpr auto & eta(int32_t i) { return eta_(i); } + constexpr auto &eta(int32_t i) { return eta_(i); } constexpr auto eta(int32_t i) const { return eta_(i); } // pt accessors - constexpr auto & pt(int32_t i) { return pt_(i); } + constexpr auto &pt(int32_t i) { return pt_(i); } constexpr auto pt(int32_t i) const { return pt_(i); } constexpr float charge(int32_t i) const { return std::copysign(1.f, stateAtBS_.state(i)(2)); } @@ -46,13 +45,13 @@ class TrackSoAHeterogeneousT { constexpr float zip(int32_t i) const { return stateAtBS_.state(i)(4); } // hitIndices accessors - constexpr auto & hitIndices() { return hitIndices_; } - constexpr auto const & hitIndices() const { return hitIndices_; } + constexpr auto &hitIndices() { return hitIndices_; } + constexpr auto const &hitIndices() const { return hitIndices_; } // detInndices accessor constexpr int nHits(int i) const { return detIndices_.size(i); } - constexpr auto & detIndices() { return detIndices_; } - constexpr auto const & detIndices() const { return detIndices_; } + constexpr auto &detIndices() { return detIndices_; } + constexpr auto const &detIndices() const { return detIndices_; } // state at the detector of the outermost hit // representation to be decided... diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc index 87220633c9a43..e207a34b15bfa 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc @@ -37,8 +37,7 @@ using edm::ParameterSet; class PixelTrackProducer : public edm::stream::EDProducer<> { public: explicit PixelTrackProducer(const edm::ParameterSet& cfg) - : theReconstruction(cfg, consumesCollector()), - htTopoToken_(esConsumes()) { + : theReconstruction(cfg, consumesCollector()), htTopoToken_(esConsumes()) { edm::LogInfo("PixelTrackProducer") << " construction..."; produces(); produces();