diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h index f74717c41e4d7..0f0f6d21e74fc 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" @@ -19,39 +19,60 @@ 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 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); } + // 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 { @@ -70,4 +91,4 @@ namespace pixelTrack { } // namespace pixelTrack -#endif // CUDADataFormats_Track_TrackHeterogeneousT_H +#endif // CUDADataFormats_Track_PixelTrackHeterogeneousT_h diff --git a/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h b/CUDADataFormats/Track/interface/TrajectoryStateSoAT.h index 64fcd573a6991..947b2af7aabb2 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" @@ -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); @@ -56,4 +56,4 @@ struct TrajectoryStateSoAT { } }; -#endif // CUDADataFormats_Track_TrajectoryStateSOAT_H +#endif // CUDADataFormats_Track_TrajectoryStateSoAT_h 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. */ diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc index 91c3a44cc8643..e207a34b15bfa 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.cc @@ -15,38 +15,60 @@ #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()) { - 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); - edm::ESHandle httopo; - es.get().get(httopo); +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 c38fd44c0d7f5..0000000000000 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducer.h +++ /dev/null @@ -1,29 +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" - -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; -}; - -#endif // RecoPixelVertexing_PixelTrackFitting_plugins_PixelTrackProducer_h 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/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); 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..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 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;