From 7979baaf28e098ffe7dfbd1e75331c7d65039314 Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 17:40:18 +0200 Subject: [PATCH 01/23] Added amplitudeError --- .../EcalRecHitSoA/interface/EcalUncalibratedRecHit.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h b/CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h index 78c909b029dc1..7497f71269089 100644 --- a/CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h +++ b/CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h @@ -21,6 +21,7 @@ namespace ecal { typename StoragePolicy::template StorageSelector::type amplitudesAll; typename StoragePolicy::template StorageSelector::type amplitude; + typename StoragePolicy::template StorageSelector::type amplitudeError; typename StoragePolicy::template StorageSelector::type chi2; typename StoragePolicy::template StorageSelector::type pedestal; typename StoragePolicy::template StorageSelector::type jitter; @@ -32,6 +33,7 @@ namespace ecal { typename std::enable_if::value, void>::type resize(size_t size) { amplitudesAll.resize(size * EcalDataFrame::MAXSAMPLES); amplitude.resize(size); + amplitudeError.resize(size); pedestal.resize(size); chi2.resize(size); did.resize(size); From 2cf6995fc0089b7fff3b5067a564517a816633d7 Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 17:53:01 +0200 Subject: [PATCH 02/23] Added GPU weights reconstruction module for Phase2EcalRecoGPU --- .../EcalUncalibRecHitPhase2WeightsAlgoGPU.cu | 35 ++++++ .../EcalUncalibRecHitPhase2WeightsAlgoGPU.h | 19 +++ .../EcalUncalibRecHitPhase2WeightsKernels.cu | 62 ++++++++++ .../EcalUncalibRecHitPhase2WeightsKernels.h | 20 +++ ...alUncalibRecHitPhase2WeightsProducerGPU.cc | 117 ++++++++++++++++++ 5 files changed, 253 insertions(+) create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu new file mode 100644 index 0000000000000..20e9849b3959f --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu @@ -0,0 +1,35 @@ +#include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" +#include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" + +#include "EcalUncalibRecHitPhase2WeightsKernels.h" +#include "EcalUncalibRecHitPhase2WeightsAlgoGPU.h" + +namespace ecal { + namespace weights { + + void entryPoint(ecal::DigisCollection const& digis, + EventOutputDataGPU& eventOutputGPU, + cms::cuda::device::unique_ptr& weights_d, + cudaStream_t cudaStream) { + unsigned int totalChannels = digis.size; + // 64 threads per block best occupancy from Nsight compute profiler + unsigned int nchannels_per_block = 64; + unsigned int threads_1d = nchannels_per_block; + unsigned int blocks_1d = (totalChannels / threads_1d) + 1; + // shared bytes from size of weight constants, digi samples per block, uncalib rechits amplitudes per block + int shared_bytes = EcalDataFrame_Ph2::MAXSAMPLES * sizeof(double) + + nchannels_per_block * (EcalDataFrame_Ph2::MAXSAMPLES * (sizeof(uint16_t)) + sizeof(float)); + Phase2WeightsKernel<<>>( + digis.data.get(), + digis.ids.get(), + eventOutputGPU.recHits.amplitude.get(), + eventOutputGPU.recHits.amplitudeError.get(), + eventOutputGPU.recHits.did.get(), + totalChannels, + weights_d.get(), + eventOutputGPU.recHits.flags.get()); + cudaCheck(cudaGetLastError()); + } + + } // namespace weights +} // namespace ecal \ No newline at end of file diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h new file mode 100644 index 0000000000000..5519879c58479 --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h @@ -0,0 +1,19 @@ +#ifndef RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsAlgoGPU_h +#define RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsAlgoGPU_h + +#include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" + +#include "DeclsForKernelsPh2.h" + +namespace ecal { + namespace weights { + + void entryPoint(ecal::DigisCollection const&, + EventOutputDataGPU&, + cms::cuda::device::unique_ptr&, + cudaStream_t); + + } // namespace weights +} // namespace ecal + +#endif // RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsAlgoGPU_h \ No newline at end of file diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu new file mode 100644 index 0000000000000..04adbfb3452cf --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu @@ -0,0 +1,62 @@ +#include + +#include "FWCore/Utilities/interface/CMSUnrollLoop.h" +#include "DataFormats/EcalRecHit/interface/EcalUncalibratedRecHit.h" +#include "DataFormats/EcalDigi/interface/EcalLiteDTUSample.h" +#include "DataFormats/EcalDigi/interface/EcalConstants.h" + +#include "EcalUncalibRecHitPhase2WeightsKernels.h" + +namespace ecal { + namespace weights { + + __global__ void Phase2WeightsKernel(uint16_t const* digis_in, + uint32_t const* dids, + ::ecal::reco::StorageScalarType* amplitude, + ::ecal::reco::StorageScalarType* amplitudeError, + uint32_t* dids_out, + int const nchannels, + double* weights, + uint32_t* flags) { + constexpr int nsamples = EcalDataFrame_Ph2::MAXSAMPLES; + int const tx = threadIdx.x + blockIdx.x * blockDim.x; + unsigned int nchannels_per_block = blockDim.x; + + if (tx < nchannels) { + auto const did = DetId{dids[tx]}; + //dynamic shared memory + extern __shared__ char shared_mem[]; + double* shr_weights = (double*)&shared_mem[0]; + float* shr_amp = (float*)&shared_mem[nsamples * sizeof(double)]; + uint16_t* shr_digis = (uint16_t*)&shared_mem[nsamples * sizeof(double)+ nchannels_per_block * sizeof(float)]; + + shr_weights = weights; + + unsigned int bx = blockIdx.x; //block index + unsigned int btx = threadIdx.x; + + for (int sample = 0; sample < nsamples; ++sample) { + const unsigned int idx = threadIdx.x * nsamples + sample; + shr_digis[idx] = digis_in[bx * nchannels_per_block * nsamples + idx]; + } + + shr_amp[btx] = 0.0; + CMS_UNROLL_LOOP + for (int sample = 0; sample < nsamples; ++sample) { + const unsigned int idx = threadIdx.x * nsamples + sample; + const auto shr_digi = shr_digis[idx]; + shr_amp[btx] += (static_cast(ecalLiteDTU::adc(shr_digi)) * ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * + shr_weights[sample]); + } + amplitude[tx] = shr_amp[btx]; + amplitudeError[tx] = 1.0f; + dids_out[tx] = did.rawId(); + flags[tx] = 0; + if (ecalLiteDTU::gainId(shr_digis[btx * nsamples + nsamples - 1])) { + flags[tx] = 0x1 << EcalUncalibratedRecHit::kHasSwitchToGain1; + } + + } //if within nchannels + } //kernel + } //namespace weights +} //namespace ecal \ No newline at end of file diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h new file mode 100644 index 0000000000000..4517b75d53d1f --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h @@ -0,0 +1,20 @@ +#ifndef RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsKernels_h +#define RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsKernels_h + +#include "DeclsForKernelsPh2.h" + +namespace ecal { + namespace weights { + + __global__ void Phase2WeightsKernel(uint16_t const* digis_in_eb, + uint32_t const* dids_eb, + ::ecal::reco::StorageScalarType* amplitudeEB, + ::ecal::reco::StorageScalarType* amplitudeErrorEB, + uint32_t* dids_outEB, + int const nchannels, + double* weights_d, + uint32_t* flagsEB); + } //namespace weights +} //namespace ecal + +#endif \ No newline at end of file diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc new file mode 100644 index 0000000000000..254bc94c3f619 --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc @@ -0,0 +1,117 @@ +#include "FWCore/Framework/interface/stream/EDProducer.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/EventSetup.h" +#include "FWCore/Framework/interface/MakerMacros.h" + +#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" + +#include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" + +#include "EcalUncalibRecHitPhase2WeightsAlgoGPU.h" +#include "DeclsForKernelsPh2.h" + +class EcalUncalibRecHitPhase2WeightsProducerGPU : public edm::stream::EDProducer { +public: + explicit EcalUncalibRecHitPhase2WeightsProducerGPU(edm::ParameterSet const &ps); + ~EcalUncalibRecHitPhase2WeightsProducerGPU() override = default; + static void fillDescriptions(edm::ConfigurationDescriptions &); + +private: + void acquire(edm::Event const &, edm::EventSetup const &, edm::WaitingTaskWithArenaHolder) override; + void produce(edm::Event &, edm::EventSetup const &) override; + +private: + const std::vector weights_; + + using InputProduct = cms::cuda::Product>; + const edm::EDGetTokenT digisToken_; + using OutputProduct = cms::cuda::Product>; + const edm::EDPutTokenT recHitsToken_; + + // event data + ecal::weights::EventOutputDataGPU eventOutputDataGPU_; + + cms::cuda::ContextState cudaState_; + + uint32_t n_; +}; + +// constructor with initialisation of elements +EcalUncalibRecHitPhase2WeightsProducerGPU::EcalUncalibRecHitPhase2WeightsProducerGPU(const edm::ParameterSet &ps) + : weights_(ps.getParameter>("weights")), + digisToken_{consumes(ps.getParameter("digisLabelEB"))}, + recHitsToken_{produces(ps.getParameter("recHitsLabelEB"))} {} + +void EcalUncalibRecHitPhase2WeightsProducerGPU::fillDescriptions(edm::ConfigurationDescriptions &descriptions) { + edm::ParameterSetDescription desc; + + desc.add("recHitsLabelEB", "EcalUncalibRecHitsEB"); + desc.add>("weights", + {-0.121016, + -0.119899, + -0.120923, + -0.0848959, + 0.261041, + 0.509881, + 0.373591, + 0.134899, + -0.0233605, + -0.0913195, + -0.112452, + -0.118596, + -0.121737, + -0.121737, + -0.121737, + -0.121737}); + + desc.add("digisLabelEB", edm::InputTag("simEcalUnsuppressedDigis", "")); + + descriptions.addWithDefaultLabel(desc); +} + +// aquire function which initislises objects on host and device to their actual objects and calls kernal +void EcalUncalibRecHitPhase2WeightsProducerGPU::acquire(edm::Event const &event, + edm::EventSetup const &setup, + edm::WaitingTaskWithArenaHolder holder) { + // cuda products + auto const &digisProduct = event.get(digisToken_); + // raii + cms::cuda::ScopedContextAcquire ctx{digisProduct, std::move(holder), cudaState_}; + + // get actual obj + auto const &digis = ctx.get(digisProduct); + + n_ = digis.size; + + // if no digis stop here + if (n_ == 0) + return; + + // weights to GPU + + cms::cuda::device::unique_ptr weights_d = + cms::cuda::make_device_unique(EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); + + cudaCheck(cudaMemcpyAsync(weights_d.get(), + weights_.data(), + EcalDataFrame_Ph2::MAXSAMPLES * sizeof(double), + cudaMemcpyHostToDevice, + ctx.stream())); + + // output on GPU + eventOutputDataGPU_.allocate(n_, ctx.stream()); + + ecal::weights::entryPoint(digis, eventOutputDataGPU_, weights_d, ctx.stream()); +} + +void EcalUncalibRecHitPhase2WeightsProducerGPU::produce(edm::Event &event, const edm::EventSetup &setup) { + cms::cuda::ScopedContextProduce ctx{cudaState_}; + + // set the size of digis + eventOutputDataGPU_.recHits.size = n_; + + // put into the event + ctx.emplace(event, recHitsToken_, std::move(eventOutputDataGPU_.recHits)); +} + +DEFINE_FWK_MODULE(EcalUncalibRecHitPhase2WeightsProducerGPU); From c9ab238ad080a2367f6a2eb7710a4f2651fb16f2 Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 17:55:13 +0200 Subject: [PATCH 03/23] Modified EcalCPUUncalibRecHitProducer to be Phase 1 & 2 compatible --- .../plugins/EcalCPUUncalibRecHitProducer.cc | 71 +++++++++++-------- 1 file changed, 43 insertions(+), 28 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc index 801d378c7c391..adb4e858457b9 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc @@ -27,6 +27,7 @@ class EcalCPUUncalibRecHitProducer : public edm::stream::EDProducer>; edm::EDGetTokenT recHitsInEBToken_, recHitsInEEToken_; using OutputProduct = ecal::UncalibratedRecHit>; @@ -40,20 +41,27 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio edm::ParameterSetDescription desc; desc.add("recHitsInLabelEB", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEB"}); - desc.add("recHitsInLabelEE", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"}); desc.add("recHitsOutLabelEB", "EcalUncalibRecHitsEB"); - desc.add("recHitsOutLabelEE", "EcalUncalibRecHitsEE"); + desc.add("containsTimingInformation", false); + desc.add("produceEE", true); + + desc.add("recHitsInLabelEE", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"}); + desc.add("recHitsOutLabelEE", "EcalUncalibRecHitsEE"); confDesc.add("ecalCPUUncalibRecHitProducer", desc); } EcalCPUUncalibRecHitProducer::EcalCPUUncalibRecHitProducer(const edm::ParameterSet& ps) - : recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, - recHitsInEEToken_{consumes(ps.getParameter("recHitsInLabelEE"))}, + : produceEE_{ps.getParameter("produceEE")}, + recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, recHitsOutEBToken_{produces(ps.getParameter("recHitsOutLabelEB"))}, - recHitsOutEEToken_{produces(ps.getParameter("recHitsOutLabelEE"))}, - containsTimingInformation_{ps.getParameter("containsTimingInformation")} {} + containsTimingInformation_{ps.getParameter("containsTimingInformation")} { + if (produceEE_) { + recHitsInEEToken_ = consumes(ps.getParameter("recHitsInLabelEE")); + recHitsOutEEToken_ = produces(ps.getParameter("recHitsOutLabelEE")); + } +} EcalCPUUncalibRecHitProducer::~EcalCPUUncalibRecHitProducer() {} @@ -62,14 +70,11 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, edm::WaitingTaskWithArenaHolder taskHolder) { // retrieve data/ctx auto const& ebRecHitsProduct = event.get(recHitsInEBToken_); - auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); cms::cuda::ScopedContextAcquire ctx{ebRecHitsProduct, std::move(taskHolder)}; auto const& ebRecHits = ctx.get(ebRecHitsProduct); - auto const& eeRecHits = ctx.get(eeRecHitsProduct); // resize the output buffers recHitsEB_.resize(ebRecHits.size); - recHitsEE_.resize(eeRecHits.size); auto lambdaToTransfer = [&ctx](auto& dest, auto* src) { using vector_type = typename std::remove_reference::type; @@ -79,42 +84,52 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream())); }; + if (produceEE_) { + auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); + auto const& eeRecHits = ctx.get(eeRecHitsProduct); + recHitsEE_.resize(eeRecHits.size); + } // enqeue transfers lambdaToTransfer(recHitsEB_.did, ebRecHits.did.get()); - lambdaToTransfer(recHitsEE_.did, eeRecHits.did.get()); - lambdaToTransfer(recHitsEB_.amplitudesAll, ebRecHits.amplitudesAll.get()); - lambdaToTransfer(recHitsEE_.amplitudesAll, eeRecHits.amplitudesAll.get()); - lambdaToTransfer(recHitsEB_.amplitude, ebRecHits.amplitude.get()); - lambdaToTransfer(recHitsEE_.amplitude, eeRecHits.amplitude.get()); - + lambdaToTransfer(recHitsEB_.amplitudeError, ebRecHits.amplitudeError.get()); lambdaToTransfer(recHitsEB_.chi2, ebRecHits.chi2.get()); - lambdaToTransfer(recHitsEE_.chi2, eeRecHits.chi2.get()); - lambdaToTransfer(recHitsEB_.pedestal, ebRecHits.pedestal.get()); - lambdaToTransfer(recHitsEE_.pedestal, eeRecHits.pedestal.get()); - lambdaToTransfer(recHitsEB_.flags, ebRecHits.flags.get()); - lambdaToTransfer(recHitsEE_.flags, eeRecHits.flags.get()); - if (containsTimingInformation_) { lambdaToTransfer(recHitsEB_.jitter, ebRecHits.jitter.get()); - lambdaToTransfer(recHitsEE_.jitter, eeRecHits.jitter.get()); - lambdaToTransfer(recHitsEB_.jitterError, ebRecHits.jitterError.get()); - lambdaToTransfer(recHitsEE_.jitterError, eeRecHits.jitterError.get()); + } + + if (produceEE_) { + auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); + auto const& eeRecHits = ctx.get(eeRecHitsProduct); + recHitsEE_.resize(eeRecHits.size); + lambdaToTransfer(recHitsEE_.did, eeRecHits.did.get()); + lambdaToTransfer(recHitsEE_.amplitudesAll, eeRecHits.amplitudesAll.get()); + lambdaToTransfer(recHitsEE_.amplitude, eeRecHits.amplitude.get()); + lambdaToTransfer(recHitsEE_.amplitudeError, eeRecHits.amplitudeError.get()); + lambdaToTransfer(recHitsEE_.chi2, eeRecHits.chi2.get()); + lambdaToTransfer(recHitsEE_.pedestal, eeRecHits.pedestal.get()); + lambdaToTransfer(recHitsEE_.flags, eeRecHits.flags.get()); + if (containsTimingInformation_) { + lambdaToTransfer(recHitsEE_.jitter, eeRecHits.jitter.get()); + lambdaToTransfer(recHitsEE_.jitterError, eeRecHits.jitterError.get()); + } } } void EcalCPUUncalibRecHitProducer::produce(edm::Event& event, edm::EventSetup const& setup) { // tmp vectors auto recHitsOutEB = std::make_unique(std::move(recHitsEB_)); - auto recHitsOutEE = std::make_unique(std::move(recHitsEE_)); - // put into event event.put(recHitsOutEBToken_, std::move(recHitsOutEB)); - event.put(recHitsOutEEToken_, std::move(recHitsOutEE)); + + if (produceEE_) { + auto recHitsOutEE = std::make_unique(std::move(recHitsEE_)); + event.put(recHitsOutEEToken_, std::move(recHitsOutEE)); + } } -DEFINE_FWK_MODULE(EcalCPUUncalibRecHitProducer); +DEFINE_FWK_MODULE(EcalCPUUncalibRecHitProducer); \ No newline at end of file From 491e68849c6556b5237fdca0f12aadaa9f77c5cb Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 17:56:56 +0200 Subject: [PATCH 04/23] Modified EcalUncalibRecHitConvertGPU2CPUFormat to be Phase 1 & 2 compatible --- .../EcalUncalibRecHitConvertGPU2CPUFormat.cc | 89 ++++++++++--------- 1 file changed, 49 insertions(+), 40 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc index b26fbe3a0c572..74ea6cc69bb8b 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc @@ -19,73 +19,82 @@ class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { void produce(edm::Event&, edm::EventSetup const&) override; private: + bool produceEE_; const edm::EDGetTokenT recHitsGPUEB_; - const edm::EDGetTokenT recHitsGPUEE_; + edm::EDGetTokenT recHitsGPUEE_; - const std::string recHitsLabelCPUEB_, recHitsLabelCPUEE_; + const std::string recHitsLabelCPUEB_; + std::string recHitsLabelCPUEE_; }; void EcalUncalibRecHitConvertGPU2CPUFormat::fillDescriptions(edm::ConfigurationDescriptions& confDesc) { edm::ParameterSetDescription desc; desc.add("recHitsLabelGPUEB", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEB")); - desc.add("recHitsLabelGPUEE", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE")); desc.add("recHitsLabelCPUEB", "EcalUncalibRecHitsEB"); + + desc.add("produceEE", true); + + desc.add("recHitsLabelGPUEE", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE")); desc.add("recHitsLabelCPUEE", "EcalUncalibRecHitsEE"); confDesc.add("ecalUncalibRecHitConvertGPU2CPUFormat", desc); } EcalUncalibRecHitConvertGPU2CPUFormat::EcalUncalibRecHitConvertGPU2CPUFormat(const edm::ParameterSet& ps) - : recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, - recHitsGPUEE_{consumes(ps.getParameter("recHitsLabelGPUEE"))}, - recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")}, - recHitsLabelCPUEE_{ps.getParameter("recHitsLabelCPUEE")} { + : produceEE_{ps.getParameter("produceEE")}, + recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, + recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")} { produces(recHitsLabelCPUEB_); - produces(recHitsLabelCPUEE_); + if (produceEE_) { + recHitsGPUEE_ = consumes(ps.getParameter("recHitsLabelGPUEE")); + recHitsLabelCPUEE_ = ps.getParameter("recHitsLabelCPUEE"); + produces(recHitsLabelCPUEE_); + } } EcalUncalibRecHitConvertGPU2CPUFormat::~EcalUncalibRecHitConvertGPU2CPUFormat() {} void EcalUncalibRecHitConvertGPU2CPUFormat::produce(edm::Event& event, edm::EventSetup const& setup) { - edm::Handle hRecHitsGPUEB, hRecHitsGPUEE; - event.getByToken(recHitsGPUEB_, hRecHitsGPUEB); - event.getByToken(recHitsGPUEE_, hRecHitsGPUEE); - + auto const& recHitsGPUEB = event.get(recHitsGPUEB_); auto recHitsCPUEB = std::make_unique(); - auto recHitsCPUEE = std::make_unique(); - recHitsCPUEB->reserve(hRecHitsGPUEB->amplitude.size()); - recHitsCPUEE->reserve(hRecHitsGPUEE->amplitude.size()); - - for (uint32_t i = 0; i < hRecHitsGPUEB->amplitude.size(); ++i) { - recHitsCPUEB->emplace_back(DetId{hRecHitsGPUEB->did[i]}, - hRecHitsGPUEB->amplitude[i], - hRecHitsGPUEB->pedestal[i], - hRecHitsGPUEB->jitter[i], - hRecHitsGPUEB->chi2[i], - hRecHitsGPUEB->flags[i]); - (*recHitsCPUEB)[i].setJitterError(hRecHitsGPUEB->jitterError[i]); + recHitsCPUEB->reserve(recHitsGPUEB.amplitude.size()); + + for (uint32_t i = 0; i < recHitsGPUEB.amplitude.size(); ++i) { + recHitsCPUEB->emplace_back(DetId{recHitsGPUEB.did[i]}, + recHitsGPUEB.amplitude[i], + recHitsGPUEB.pedestal[i], + recHitsGPUEB.jitter[i], + recHitsGPUEB.chi2[i], + recHitsGPUEB.flags[i]); + (*recHitsCPUEB)[i].setAmplitudeError(recHitsGPUEB.amplitudeError[i]); + (*recHitsCPUEB)[i].setJitterError(recHitsGPUEB.jitterError[i]); auto const offset = i * EcalDataFrame::MAXSAMPLES; for (uint32_t sample = 0; sample < EcalDataFrame::MAXSAMPLES; ++sample) - (*recHitsCPUEB)[i].setOutOfTimeAmplitude(sample, hRecHitsGPUEB->amplitudesAll[offset + sample]); + (*recHitsCPUEB)[i].setOutOfTimeAmplitude(sample, recHitsGPUEB.amplitudesAll[offset + sample]); } - - for (uint32_t i = 0; i < hRecHitsGPUEE->amplitude.size(); ++i) { - recHitsCPUEE->emplace_back(DetId{hRecHitsGPUEE->did[i]}, - hRecHitsGPUEE->amplitude[i], - hRecHitsGPUEE->pedestal[i], - hRecHitsGPUEE->jitter[i], - hRecHitsGPUEE->chi2[i], - hRecHitsGPUEE->flags[i]); - (*recHitsCPUEE)[i].setJitterError(hRecHitsGPUEE->jitterError[i]); - auto const offset = i * EcalDataFrame::MAXSAMPLES; - for (uint32_t sample = 0; sample < EcalDataFrame::MAXSAMPLES; ++sample) - (*recHitsCPUEE)[i].setOutOfTimeAmplitude(sample, hRecHitsGPUEE->amplitudesAll[offset + sample]); + if (produceEE_) { + auto const& recHitsGPUEE = event.get(recHitsGPUEE_); + auto recHitsCPUEE = std::make_unique(); + recHitsCPUEE->reserve(recHitsGPUEE.amplitude.size()); + for (uint32_t i = 0; i < recHitsGPUEE.amplitude.size(); ++i) { + recHitsCPUEE->emplace_back(DetId{recHitsGPUEE.did[i]}, + recHitsGPUEE.amplitude[i], + recHitsGPUEE.pedestal[i], + recHitsGPUEE.jitter[i], + recHitsGPUEE.chi2[i], + recHitsGPUEE.flags[i]); + (*recHitsCPUEB)[i].setAmplitudeError(recHitsGPUEE.amplitudeError[i]); + (*recHitsCPUEE)[i].setJitterError(recHitsGPUEE.jitterError[i]); + auto const offset = i * EcalDataFrame::MAXSAMPLES; + for (uint32_t sample = 0; sample < EcalDataFrame::MAXSAMPLES; ++sample) { + (*recHitsCPUEE)[i].setOutOfTimeAmplitude(sample, recHitsGPUEE.amplitudesAll[offset + sample]); + } + } + event.put(std::move(recHitsCPUEE), recHitsLabelCPUEE_); } - event.put(std::move(recHitsCPUEB), recHitsLabelCPUEB_); - event.put(std::move(recHitsCPUEE), recHitsLabelCPUEE_); } -DEFINE_FWK_MODULE(EcalUncalibRecHitConvertGPU2CPUFormat); +DEFINE_FWK_MODULE(EcalUncalibRecHitConvertGPU2CPUFormat); \ No newline at end of file From 5fd429d59d0ceb72002f806f76bd1f1a527be9a3 Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 17:59:07 +0200 Subject: [PATCH 05/23] Added configuration for testing Phase2EcalRecoGPU --- .../EcalRecProducers/python/ecalUncalibRecHitPhase2GPU_cfi.py | 2 ++ 1 file changed, 2 insertions(+) create mode 100644 RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2GPU_cfi.py diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2GPU_cfi.py b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2GPU_cfi.py new file mode 100644 index 0000000000000..b78732e554bca --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2GPU_cfi.py @@ -0,0 +1,2 @@ +import RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2WeightsProducerGPU_cfi as _mod +ecalUncalibRecHitPhase2GPU = _mod.ecalUncalibRecHitPhase2WeightsProducerGPU.clone() \ No newline at end of file From e35ae4dfc405ff37c66ae77d0c6fdf37b539855c Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 18:00:31 +0200 Subject: [PATCH 06/23] Modified ecalUncalibRecHitPhase2_cff to use a switch producer --- .../python/ecalUncalibRecHitPhase2_cff.py | 50 ++++++++++++++++++- 1 file changed, 49 insertions(+), 1 deletion(-) diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py index 63a3f7711dde3..2ac18fd0a9d09 100644 --- a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py +++ b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py @@ -1,4 +1,52 @@ import FWCore.ParameterSet.Config as cms +from HeterogeneousCore.CUDACore.SwitchProducerCUDA import SwitchProducerCUDA +from Configuration.ProcessModifiers.gpu_cff import gpu -from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2_cfi import * + +from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2_cfi import ecalUncalibRecHitPhase2 as _ecalUncalibRecHitPhase2 +ecalUncalibRecHitPhase2 = SwitchProducerCUDA( + cpu = _ecalUncalibRecHitPhase2.clone() +) + +# cpu weights ecalUncalibRecHitPhase2Task = cms.Task(ecalUncalibRecHitPhase2) + +# conditions used on gpu + + +from RecoLocalCalo.EcalRecProducers.ecalPh2DigiToGPUProducer_cfi import ecalPh2DigiToGPUProducer as _ecalPh2DigiToGPUProducer +ecalPh2DigiToGPUProducer = _ecalPh2DigiToGPUProducer.clone() + +# gpu weights +from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2GPU_cfi import ecalUncalibRecHitPhase2GPU as _ecalUncalibRecHitPhase2GPU +ecalUncalibRecHitPhase2GPU = _ecalUncalibRecHitPhase2GPU.clone( + digisLabelEB = cms.InputTag('ecalPh2DigiToGPUProducer', 'ebDigis') +) + +# copy the uncalibrated rechits from GPU to CPU +from RecoLocalCalo.EcalRecProducers.ecalCPUUncalibRecHitProducer_cfi import ecalCPUUncalibRecHitProducer as _ecalCPUUncalibRecHitProducer +ecalMultiFitUncalibRecHitSoAnew = _ecalCPUUncalibRecHitProducer.clone( + recHitsInLabelEB = cms.InputTag('ecalUncalibRecHitPhase2GPU', 'EcalUncalibRecHitsEB'), + produceEE = cms.bool(False) +) + + +# convert the uncalibrated rechits from SoA to legacy format +from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitConvertGPU2CPUFormat_cfi import ecalUncalibRecHitConvertGPU2CPUFormat as _ecalUncalibRecHitConvertGPU2CPUFormat +gpu.toModify(ecalUncalibRecHitPhase2, + cuda = _ecalUncalibRecHitConvertGPU2CPUFormat.clone( + recHitsLabelGPUEB = cms.InputTag('ecalMultiFitUncalibRecHitSoAnew', 'EcalUncalibRecHitsEB'), + produceEE = cms.bool(False) + ) +) + +gpu.toReplaceWith(ecalUncalibRecHitPhase2Task, cms.Task( + # convert phase2 digis to GPU SoA + ecalPh2DigiToGPUProducer, + # ECAL weights running on GPU + ecalUncalibRecHitPhase2GPU, + # copy the uncalibrated rechits from GPU to CPU + ecalMultiFitUncalibRecHitSoAnew, + # ECAL multifit running on CPU, or convert the uncalibrated rechits from SoA to legacy format + ecalUncalibRecHitPhase2, +)) From b6eb69a51ae64429fe25711167290f86d7258fd8 Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 26 Apr 2022 18:02:39 +0200 Subject: [PATCH 07/23] Added test configuration for Phase2EcalRecoGPU --- ...alUncalibRecHitPhase2WeightsProducerGPU.py | 142 ++++++++++++++++++ 1 file changed, 142 insertions(+) create mode 100644 RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU.py diff --git a/RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU.py b/RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU.py new file mode 100644 index 0000000000000..93e02b619917a --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU.py @@ -0,0 +1,142 @@ +import FWCore.ParameterSet.Config as cms + +from Configuration.Eras.Era_Phase2C10_cff import Phase2C10 +from Configuration.Eras.Modifier_phase2_ecal_devel_cff import phase2_ecal_devel +from Configuration.ProcessModifiers.gpu_cff import gpu +from Configuration.ProcessModifiers.gpuValidationEcal_cff import gpuValidationEcal + +process = cms.Process('RECO',Phase2C10,phase2_ecal_devel, gpu, gpuValidationEcal) + +# import of standard configurations +process.load('Configuration.StandardSequences.Services_cff') +process.load('SimGeneral.HepPDTESSource.pythiapdt_cfi') +process.load('FWCore.MessageService.MessageLogger_cfi') +process.load('Configuration.EventContent.EventContent_cff') +process.load('SimGeneral.MixingModule.mixNoPU_cfi') +process.load('Configuration.Geometry.GeometryExtended2026D60Reco_cff') +process.load('Configuration.StandardSequences.MagneticField_cff') +process.load('Configuration.StandardSequences.RawToDigi_cff') +process.load('Configuration.StandardSequences.L1Reco_cff') +process.load('Configuration.StandardSequences.Reconstruction_cff') +process.load('Configuration.StandardSequences.FrontierConditions_GlobalTag_cff') +process.load('DQMOffline.Configuration.DQMOfflineMC_cff') +process.load('RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2_cff') +process.load('RecoLuminosity.LumiProducer.bunchSpacingProducer_cfi') + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32(100), + output = cms.optional.untracked.allowed(cms.int32,cms.PSet) +) + +process.load('HLTrigger.Timer.FastTimerService_cfi') +process.FastTimerService.enableDQM = False +process.FastTimerService.printRunSummary = False +process.FastTimerService.printJobSummary = True +process.FastTimerService.writeJSONSummary = True +process.FastTimerService.jsonFileName = 'resources.json' +process.MessageLogger.FastReport = cms.untracked.PSet() + + + +# Input source +process.source = cms.Source("PoolSource", + fileNames = cms.untracked.vstring('/store/group/dpg_ecal/comm_ecal/upgrade/Phase2CMSSW///RelValTTbar_14TeV_ecaldigi_123X_mcRun4_realistic_v3_2026D77noPU-v1.root'), + secondaryFileNames = cms.untracked.vstring() +) + +process.options = cms.untracked.PSet( + FailPath = cms.untracked.vstring(), + IgnoreCompletely = cms.untracked.vstring(), + Rethrow = cms.untracked.vstring(), + SkipEvent = cms.untracked.vstring(), + allowUnscheduled = cms.obsolete.untracked.bool, + canDeleteEarly = cms.untracked.vstring(), + emptyRunLumiMode = cms.obsolete.untracked.string, + eventSetup = cms.untracked.PSet( + forceNumberOfConcurrentIOVs = cms.untracked.PSet( + allowAnyLabel_=cms.required.untracked.uint32 + ), + numberOfConcurrentIOVs = cms.untracked.uint32(1) + ), + fileMode = cms.untracked.string('FULLMERGE'), + forceEventSetupCacheClearOnNewRun = cms.untracked.bool(False), + makeTriggerResults = cms.obsolete.untracked.bool, + numberOfConcurrentLuminosityBlocks = cms.untracked.uint32(1), + numberOfConcurrentRuns = cms.untracked.uint32(1), + numberOfStreams = cms.untracked.uint32(0), + numberOfThreads = cms.untracked.uint32(1), + printDependencies = cms.untracked.bool(False), + sizeOfStackForThreadsInKB = cms.optional.untracked.uint32, + throwIfIllegalParameter = cms.untracked.bool(True), + wantSummary = cms.untracked.bool(True) +) + +# Production Info +process.configurationMetadata = cms.untracked.PSet( + annotation = cms.untracked.string('step3GPU nevts:10'), + name = cms.untracked.string('Applications'), + version = cms.untracked.string('$Revision: 1.19 $') +) + +# Set up the DQM GPU validation task +process.ecalMonitorTaskEcalOnly.workers = ["GpuTask"] +process.ecalMonitorTaskEcalOnly.collectionTags.EBCpuUncalibRecHit = "ecalUncalibRecHitPhase2@cpu:EcalUncalibRecHitsEB" +process.ecalMonitorTaskEcalOnly.collectionTags.EBGpuUncalibRecHit = "ecalUncalibRecHitPhase2@cuda:EcalUncalibRecHitsEB" + +# Output definition +outputCommand = process.FEVTDEBUGHLTEventContent.outputCommands +outputCommand.append('keep *_ecalUncalibRecHitPhase2*_*_RECO') +process.FEVTDEBUGHLToutput = cms.OutputModule("PoolOutputModule", + dataset = cms.untracked.PSet( + dataTier = cms.untracked.string('GEN-SIM-RECO'), + filterName = cms.untracked.string('') + ), + fileName = cms.untracked.string('file:step3GPU.root'), + outputCommands = outputCommand, + splitLevel = cms.untracked.int32(0) +) + +process.DQMoutput = cms.OutputModule("DQMRootOutputModule", + dataset = cms.untracked.PSet( + dataTier = cms.untracked.string('DQMIO'), + filterName = cms.untracked.string('') + ), + fileName = cms.untracked.string('file:step3GPU_inDQM.root'), + outputCommands = process.DQMEventContent.outputCommands, + splitLevel = cms.untracked.int32(0) +) + +# Additional output definition + +# Other statements +from Configuration.AlCa.GlobalTag import GlobalTag +process.GlobalTag = GlobalTag(process.GlobalTag, 'auto:phase2_realistic_T15', '') + +# Path and EndPath definitions +process.L1Reco_step = cms.Path(process.L1Reco) +process.reconstruction_step = cms.Path(cms.Sequence(cms.Task( + process.bunchSpacingProducer, + process.ecalUncalibRecHitPhase2Task +))) +process.dqmoffline_step = cms.EndPath(process.DQMOfflineEcalOnly) + +process.FEVTDEBUGHLToutput_step = cms.EndPath(process.FEVTDEBUGHLToutput) +process.DQMoutput_step = cms.EndPath(process.DQMoutput) + +# Schedule definition +process.schedule = cms.Schedule(process.L1Reco_step,process.reconstruction_step,process.dqmoffline_step,process.FEVTDEBUGHLToutput_step,process.DQMoutput_step) +from PhysicsTools.PatAlgos.tools.helpers import associatePatAlgosToolsTask +associatePatAlgosToolsTask(process) + + + +# Customisation from command line + +#Have logErrorHarvester wait for the same EDProducers to finish as those providing data for the OutputModule +from FWCore.Modules.logErrorHarvester_cff import customiseLogErrorHarvesterUsingOutputCommands +process = customiseLogErrorHarvesterUsingOutputCommands(process) + +# Add early deletion of temporary data products to reduce peak memory need +from Configuration.StandardSequences.earlyDeleteSettings_cff import customiseEarlyDelete +process = customiseEarlyDelete(process) +# End adding early deletion \ No newline at end of file From 724935fa6fa9bfe219073ee1c4a8bc0e9d9072ed Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Fri, 29 Apr 2022 12:38:51 +0200 Subject: [PATCH 08/23] Added isPhase1 for amplitudeError issue --- .../plugins/EcalCPUUncalibRecHitProducer.cc | 18 +++++++++--------- .../EcalUncalibRecHitConvertGPU2CPUFormat.cc | 14 +++++++------- .../python/ecalUncalibRecHitPhase2_cff.py | 4 ++-- 3 files changed, 18 insertions(+), 18 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc index adb4e858457b9..2b87a3dd565d8 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc @@ -27,7 +27,7 @@ class EcalCPUUncalibRecHitProducer : public edm::stream::EDProducer>; edm::EDGetTokenT recHitsInEBToken_, recHitsInEEToken_; using OutputProduct = ecal::UncalibratedRecHit>; @@ -44,7 +44,7 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio desc.add("recHitsOutLabelEB", "EcalUncalibRecHitsEB"); desc.add("containsTimingInformation", false); - desc.add("produceEE", true); + desc.add("isPhase1", true); desc.add("recHitsInLabelEE", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"}); desc.add("recHitsOutLabelEE", "EcalUncalibRecHitsEE"); @@ -53,11 +53,11 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio } EcalCPUUncalibRecHitProducer::EcalCPUUncalibRecHitProducer(const edm::ParameterSet& ps) - : produceEE_{ps.getParameter("produceEE")}, + : isPhase1_{ps.getParameter("isPhase1")}, recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, recHitsOutEBToken_{produces(ps.getParameter("recHitsOutLabelEB"))}, containsTimingInformation_{ps.getParameter("containsTimingInformation")} { - if (produceEE_) { + if (isPhase1_) { recHitsInEEToken_ = consumes(ps.getParameter("recHitsInLabelEE")); recHitsOutEEToken_ = produces(ps.getParameter("recHitsOutLabelEE")); } @@ -84,7 +84,7 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream())); }; - if (produceEE_) { + if (isPhase1_) { auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); auto const& eeRecHits = ctx.get(eeRecHitsProduct); recHitsEE_.resize(eeRecHits.size); @@ -93,7 +93,6 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, lambdaToTransfer(recHitsEB_.did, ebRecHits.did.get()); lambdaToTransfer(recHitsEB_.amplitudesAll, ebRecHits.amplitudesAll.get()); lambdaToTransfer(recHitsEB_.amplitude, ebRecHits.amplitude.get()); - lambdaToTransfer(recHitsEB_.amplitudeError, ebRecHits.amplitudeError.get()); lambdaToTransfer(recHitsEB_.chi2, ebRecHits.chi2.get()); lambdaToTransfer(recHitsEB_.pedestal, ebRecHits.pedestal.get()); lambdaToTransfer(recHitsEB_.flags, ebRecHits.flags.get()); @@ -101,15 +100,16 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, lambdaToTransfer(recHitsEB_.jitter, ebRecHits.jitter.get()); lambdaToTransfer(recHitsEB_.jitterError, ebRecHits.jitterError.get()); } + if (!(isPhase1_)) + lambdaToTransfer(recHitsEB_.amplitudeError, ebRecHits.amplitudeError.get()); - if (produceEE_) { + if (isPhase1_) { auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); auto const& eeRecHits = ctx.get(eeRecHitsProduct); recHitsEE_.resize(eeRecHits.size); lambdaToTransfer(recHitsEE_.did, eeRecHits.did.get()); lambdaToTransfer(recHitsEE_.amplitudesAll, eeRecHits.amplitudesAll.get()); lambdaToTransfer(recHitsEE_.amplitude, eeRecHits.amplitude.get()); - lambdaToTransfer(recHitsEE_.amplitudeError, eeRecHits.amplitudeError.get()); lambdaToTransfer(recHitsEE_.chi2, eeRecHits.chi2.get()); lambdaToTransfer(recHitsEE_.pedestal, eeRecHits.pedestal.get()); lambdaToTransfer(recHitsEE_.flags, eeRecHits.flags.get()); @@ -126,7 +126,7 @@ void EcalCPUUncalibRecHitProducer::produce(edm::Event& event, edm::EventSetup co // put into event event.put(recHitsOutEBToken_, std::move(recHitsOutEB)); - if (produceEE_) { + if (isPhase1_) { auto recHitsOutEE = std::make_unique(std::move(recHitsEE_)); event.put(recHitsOutEEToken_, std::move(recHitsOutEE)); } diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc index 74ea6cc69bb8b..f74b1cabe8d77 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc @@ -19,7 +19,7 @@ class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { void produce(edm::Event&, edm::EventSetup const&) override; private: - bool produceEE_; + bool isPhase1_; const edm::EDGetTokenT recHitsGPUEB_; edm::EDGetTokenT recHitsGPUEE_; @@ -34,7 +34,7 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::fillDescriptions(edm::ConfigurationD desc.add("recHitsLabelCPUEB", "EcalUncalibRecHitsEB"); - desc.add("produceEE", true); + desc.add("isPhase1", true); desc.add("recHitsLabelGPUEE", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE")); desc.add("recHitsLabelCPUEE", "EcalUncalibRecHitsEE"); @@ -43,11 +43,11 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::fillDescriptions(edm::ConfigurationD } EcalUncalibRecHitConvertGPU2CPUFormat::EcalUncalibRecHitConvertGPU2CPUFormat(const edm::ParameterSet& ps) - : produceEE_{ps.getParameter("produceEE")}, + : isPhase1_{ps.getParameter("isPhase1")}, recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")} { produces(recHitsLabelCPUEB_); - if (produceEE_) { + if (isPhase1_) { recHitsGPUEE_ = consumes(ps.getParameter("recHitsLabelGPUEE")); recHitsLabelCPUEE_ = ps.getParameter("recHitsLabelCPUEE"); produces(recHitsLabelCPUEE_); @@ -68,13 +68,14 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::produce(edm::Event& event, edm::Even recHitsGPUEB.jitter[i], recHitsGPUEB.chi2[i], recHitsGPUEB.flags[i]); - (*recHitsCPUEB)[i].setAmplitudeError(recHitsGPUEB.amplitudeError[i]); + if (!(isPhase1_)) + (*recHitsCPUEB)[i].setAmplitudeError(recHitsGPUEB.amplitudeError[i]); (*recHitsCPUEB)[i].setJitterError(recHitsGPUEB.jitterError[i]); auto const offset = i * EcalDataFrame::MAXSAMPLES; for (uint32_t sample = 0; sample < EcalDataFrame::MAXSAMPLES; ++sample) (*recHitsCPUEB)[i].setOutOfTimeAmplitude(sample, recHitsGPUEB.amplitudesAll[offset + sample]); } - if (produceEE_) { + if (isPhase1_) { auto const& recHitsGPUEE = event.get(recHitsGPUEE_); auto recHitsCPUEE = std::make_unique(); recHitsCPUEE->reserve(recHitsGPUEE.amplitude.size()); @@ -85,7 +86,6 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::produce(edm::Event& event, edm::Even recHitsGPUEE.jitter[i], recHitsGPUEE.chi2[i], recHitsGPUEE.flags[i]); - (*recHitsCPUEB)[i].setAmplitudeError(recHitsGPUEE.amplitudeError[i]); (*recHitsCPUEE)[i].setJitterError(recHitsGPUEE.jitterError[i]); auto const offset = i * EcalDataFrame::MAXSAMPLES; for (uint32_t sample = 0; sample < EcalDataFrame::MAXSAMPLES; ++sample) { diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py index 2ac18fd0a9d09..6349a6b7e082e 100644 --- a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py +++ b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py @@ -27,7 +27,7 @@ from RecoLocalCalo.EcalRecProducers.ecalCPUUncalibRecHitProducer_cfi import ecalCPUUncalibRecHitProducer as _ecalCPUUncalibRecHitProducer ecalMultiFitUncalibRecHitSoAnew = _ecalCPUUncalibRecHitProducer.clone( recHitsInLabelEB = cms.InputTag('ecalUncalibRecHitPhase2GPU', 'EcalUncalibRecHitsEB'), - produceEE = cms.bool(False) + isPhase1 = cms.bool(False) ) @@ -36,7 +36,7 @@ gpu.toModify(ecalUncalibRecHitPhase2, cuda = _ecalUncalibRecHitConvertGPU2CPUFormat.clone( recHitsLabelGPUEB = cms.InputTag('ecalMultiFitUncalibRecHitSoAnew', 'EcalUncalibRecHitsEB'), - produceEE = cms.bool(False) + isPhase1 = cms.bool(False) ) ) From d8c9212d6b54628de689b0df515f3591203bbf0f Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Fri, 29 Apr 2022 12:45:24 +0200 Subject: [PATCH 09/23] re-added testEcalUncalibRecHitPhase2WeightsProducerGPU_harvesting.py --- ...cHitPhase2WeightsProducerGPU_harvesting.py | 96 +++++++++++++++++++ 1 file changed, 96 insertions(+) create mode 100644 RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU_harvesting.py diff --git a/RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU_harvesting.py b/RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU_harvesting.py new file mode 100644 index 0000000000000..e3c518b073468 --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/test/testEcalUncalibRecHitPhase2WeightsProducerGPU_harvesting.py @@ -0,0 +1,96 @@ +# Auto generated configuration file +# using: +# Revision: 1.19 +# Source: /local/reps/CMSSW/CMSSW/Configuration/Applications/python/ConfigBuilder.py,v +# with command line options: step4 -s HARVESTING:@ecalOnlyValidation+@ecal --conditions auto:phase1_2021_realistic --mc --geometry DB:Extended --scenario pp --filetype DQM --era Run3 -n 100 --filein file:step3GPU_inDQM.root --fileout file:step4.root +import FWCore.ParameterSet.Config as cms + +from Configuration.Eras.Era_Run3_cff import Run3 + +process = cms.Process('HARVESTING',Run3) + +# import of standard configurations +process.load('Configuration.StandardSequences.Services_cff') +process.load('SimGeneral.HepPDTESSource.pythiapdt_cfi') +process.load('FWCore.MessageService.MessageLogger_cfi') +process.load('Configuration.EventContent.EventContent_cff') +process.load('SimGeneral.MixingModule.mixNoPU_cfi') +process.load('Configuration.StandardSequences.GeometryRecoDB_cff') +process.load('Configuration.StandardSequences.MagneticField_cff') +process.load('Configuration.StandardSequences.DQMSaverAtRunEnd_cff') +process.load('Configuration.StandardSequences.Harvesting_cff') +process.load('Configuration.StandardSequences.FrontierConditions_GlobalTag_cff') + +process.maxEvents = cms.untracked.PSet( + input = cms.untracked.int32(100), + output = cms.optional.untracked.allowed(cms.int32,cms.PSet) +) + +# Input source +process.source = cms.Source("DQMRootSource", + fileNames = cms.untracked.vstring('file:step3GPU_inDQM.root') +) + +process.options = cms.untracked.PSet( + FailPath = cms.untracked.vstring(), + IgnoreCompletely = cms.untracked.vstring(), + Rethrow = cms.untracked.vstring('ProductNotFound'), + SkipEvent = cms.untracked.vstring(), + allowUnscheduled = cms.obsolete.untracked.bool, + canDeleteEarly = cms.untracked.vstring(), + deleteNonConsumedUnscheduledModules = cms.untracked.bool(True), + dumpOptions = cms.untracked.bool(False), + emptyRunLumiMode = cms.obsolete.untracked.string, + eventSetup = cms.untracked.PSet( + forceNumberOfConcurrentIOVs = cms.untracked.PSet( + allowAnyLabel_=cms.required.untracked.uint32 + ), + numberOfConcurrentIOVs = cms.untracked.uint32(0) + ), + fileMode = cms.untracked.string('FULLMERGE'), + forceEventSetupCacheClearOnNewRun = cms.untracked.bool(False), + makeTriggerResults = cms.obsolete.untracked.bool, + numberOfConcurrentLuminosityBlocks = cms.untracked.uint32(0), + numberOfConcurrentRuns = cms.untracked.uint32(1), + numberOfStreams = cms.untracked.uint32(0), + numberOfThreads = cms.untracked.uint32(1), + printDependencies = cms.untracked.bool(False), + sizeOfStackForThreadsInKB = cms.optional.untracked.uint32, + throwIfIllegalParameter = cms.untracked.bool(True), + wantSummary = cms.untracked.bool(False) +) + +# Production Info +process.configurationMetadata = cms.untracked.PSet( + annotation = cms.untracked.string('step4 nevts:100'), + name = cms.untracked.string('Applications'), + version = cms.untracked.string('$Revision: 1.19 $') +) + +# Output definition + +# Additional output definition + +# Other statements +from Configuration.AlCa.GlobalTag import GlobalTag +process.GlobalTag = GlobalTag(process.GlobalTag, 'auto:phase2_realistic_T15', '') + +# Path and EndPath definitions +process.postValidation_ECAL_step = cms.Path(process.postValidation_ECAL) +process.DQMHarvestEcal_step = cms.Path(process.DQMHarvestEcal) +process.DQMCertEcal_step = cms.Path(process.DQMCertEcal) +process.dqmsave_step = cms.Path(process.DQMSaver) + +# Schedule definition +process.schedule = cms.Schedule(process.postValidation_ECAL_step,process.DQMHarvestEcal_step,process.DQMCertEcal_step,process.dqmsave_step) +from PhysicsTools.PatAlgos.tools.helpers import associatePatAlgosToolsTask +associatePatAlgosToolsTask(process) + + + +# Customisation from command line + +# Add early deletion of temporary data products to reduce peak memory need +from Configuration.StandardSequences.earlyDeleteSettings_cff import customiseEarlyDelete +process = customiseEarlyDelete(process) +# End adding early deletion \ No newline at end of file From e9f49d8a64dadb3e4d9afa8a96456a26320f72fc Mon Sep 17 00:00:00 2001 From: Christopher Sandever Date: Tue, 3 May 2022 16:09:43 +0200 Subject: [PATCH 10/23] Changed isPhase1 to not(isPhase2) --- .../plugins/EcalCPUUncalibRecHitProducer.cc | 16 ++++++++-------- .../EcalUncalibRecHitConvertGPU2CPUFormat.cc | 12 ++++++------ .../python/ecalUncalibRecHitPhase2_cff.py | 4 ++-- 3 files changed, 16 insertions(+), 16 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc index 2b87a3dd565d8..cc29e24e8e9e9 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc @@ -27,7 +27,7 @@ class EcalCPUUncalibRecHitProducer : public edm::stream::EDProducer>; edm::EDGetTokenT recHitsInEBToken_, recHitsInEEToken_; using OutputProduct = ecal::UncalibratedRecHit>; @@ -44,7 +44,7 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio desc.add("recHitsOutLabelEB", "EcalUncalibRecHitsEB"); desc.add("containsTimingInformation", false); - desc.add("isPhase1", true); + desc.add("isPhase2", false); desc.add("recHitsInLabelEE", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"}); desc.add("recHitsOutLabelEE", "EcalUncalibRecHitsEE"); @@ -53,11 +53,11 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio } EcalCPUUncalibRecHitProducer::EcalCPUUncalibRecHitProducer(const edm::ParameterSet& ps) - : isPhase1_{ps.getParameter("isPhase1")}, + : isPhase2_{ps.getParameter("isPhase2")}, recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, recHitsOutEBToken_{produces(ps.getParameter("recHitsOutLabelEB"))}, containsTimingInformation_{ps.getParameter("containsTimingInformation")} { - if (isPhase1_) { + if (!isPhase2_) { recHitsInEEToken_ = consumes(ps.getParameter("recHitsInLabelEE")); recHitsOutEEToken_ = produces(ps.getParameter("recHitsOutLabelEE")); } @@ -84,7 +84,7 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream())); }; - if (isPhase1_) { + if (!isPhase2_) { auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); auto const& eeRecHits = ctx.get(eeRecHitsProduct); recHitsEE_.resize(eeRecHits.size); @@ -100,10 +100,10 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, lambdaToTransfer(recHitsEB_.jitter, ebRecHits.jitter.get()); lambdaToTransfer(recHitsEB_.jitterError, ebRecHits.jitterError.get()); } - if (!(isPhase1_)) + if (isPhase2_) lambdaToTransfer(recHitsEB_.amplitudeError, ebRecHits.amplitudeError.get()); - if (isPhase1_) { + if (!isPhase2_) { auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); auto const& eeRecHits = ctx.get(eeRecHitsProduct); recHitsEE_.resize(eeRecHits.size); @@ -126,7 +126,7 @@ void EcalCPUUncalibRecHitProducer::produce(edm::Event& event, edm::EventSetup co // put into event event.put(recHitsOutEBToken_, std::move(recHitsOutEB)); - if (isPhase1_) { + if (!isPhase2_) { auto recHitsOutEE = std::make_unique(std::move(recHitsEE_)); event.put(recHitsOutEEToken_, std::move(recHitsOutEE)); } diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc index f74b1cabe8d77..941da16f26717 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc @@ -19,7 +19,7 @@ class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { void produce(edm::Event&, edm::EventSetup const&) override; private: - bool isPhase1_; + bool isPhase2_; const edm::EDGetTokenT recHitsGPUEB_; edm::EDGetTokenT recHitsGPUEE_; @@ -34,7 +34,7 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::fillDescriptions(edm::ConfigurationD desc.add("recHitsLabelCPUEB", "EcalUncalibRecHitsEB"); - desc.add("isPhase1", true); + desc.add("isPhase2", false); desc.add("recHitsLabelGPUEE", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE")); desc.add("recHitsLabelCPUEE", "EcalUncalibRecHitsEE"); @@ -43,11 +43,11 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::fillDescriptions(edm::ConfigurationD } EcalUncalibRecHitConvertGPU2CPUFormat::EcalUncalibRecHitConvertGPU2CPUFormat(const edm::ParameterSet& ps) - : isPhase1_{ps.getParameter("isPhase1")}, + : isPhase2_{ps.getParameter("isPhase2")}, recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")} { produces(recHitsLabelCPUEB_); - if (isPhase1_) { + if (!isPhase2_) { recHitsGPUEE_ = consumes(ps.getParameter("recHitsLabelGPUEE")); recHitsLabelCPUEE_ = ps.getParameter("recHitsLabelCPUEE"); produces(recHitsLabelCPUEE_); @@ -68,14 +68,14 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::produce(edm::Event& event, edm::Even recHitsGPUEB.jitter[i], recHitsGPUEB.chi2[i], recHitsGPUEB.flags[i]); - if (!(isPhase1_)) + if (isPhase2_) (*recHitsCPUEB)[i].setAmplitudeError(recHitsGPUEB.amplitudeError[i]); (*recHitsCPUEB)[i].setJitterError(recHitsGPUEB.jitterError[i]); auto const offset = i * EcalDataFrame::MAXSAMPLES; for (uint32_t sample = 0; sample < EcalDataFrame::MAXSAMPLES; ++sample) (*recHitsCPUEB)[i].setOutOfTimeAmplitude(sample, recHitsGPUEB.amplitudesAll[offset + sample]); } - if (isPhase1_) { + if (!isPhase2_) { auto const& recHitsGPUEE = event.get(recHitsGPUEE_); auto recHitsCPUEE = std::make_unique(); recHitsCPUEE->reserve(recHitsGPUEE.amplitude.size()); diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py index 6349a6b7e082e..7b03dbf9bee92 100644 --- a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py +++ b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py @@ -27,7 +27,7 @@ from RecoLocalCalo.EcalRecProducers.ecalCPUUncalibRecHitProducer_cfi import ecalCPUUncalibRecHitProducer as _ecalCPUUncalibRecHitProducer ecalMultiFitUncalibRecHitSoAnew = _ecalCPUUncalibRecHitProducer.clone( recHitsInLabelEB = cms.InputTag('ecalUncalibRecHitPhase2GPU', 'EcalUncalibRecHitsEB'), - isPhase1 = cms.bool(False) + isPhase2 = cms.bool(True) ) @@ -36,7 +36,7 @@ gpu.toModify(ecalUncalibRecHitPhase2, cuda = _ecalUncalibRecHitConvertGPU2CPUFormat.clone( recHitsLabelGPUEB = cms.InputTag('ecalMultiFitUncalibRecHitSoAnew', 'EcalUncalibRecHitsEB'), - isPhase1 = cms.bool(False) + isPhase2 = cms.bool(True) ) ) From 5a3e7202851ccb5afd28385e6ad94ecda39b1544 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:12:37 +0100 Subject: [PATCH 11/23] Added DeclsForKernelsPhase2.h --- .../plugins/DeclsForKernelsPhase2.h | 28 +++++++++++++++++++ 1 file changed, 28 insertions(+) create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernelsPhase2.h diff --git a/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernelsPhase2.h b/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernelsPhase2.h new file mode 100644 index 0000000000000..6e52962706ee4 --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernelsPhase2.h @@ -0,0 +1,28 @@ +#ifndef RecoLocalCalo_EcalRecProducers_plugins_DeclsForKernelsPhase2_h +#define RecoLocalCalo_EcalRecProducers_plugins_DeclsForKernelsPhase2_h + +#include "CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h" +#include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" + +namespace ecal { + namespace weights { + + struct EventOutputDataGPU { + UncalibratedRecHit<::calo::common::DevStoragePolicy> recHits; + + void allocate(uint32_t digi_size, cudaStream_t cudaStream) { + auto const size = digi_size; + recHits.amplitudesAll = + cms::cuda::make_device_unique(size * EcalDataFrame::MAXSAMPLES, cudaStream); + recHits.amplitude = cms::cuda::make_device_unique(size, cudaStream); + recHits.amplitudeError = cms::cuda::make_device_unique(size, cudaStream); + recHits.chi2 = cms::cuda::make_device_unique(size, cudaStream); + recHits.pedestal = cms::cuda::make_device_unique(size, cudaStream); + recHits.did = cms::cuda::make_device_unique(size, cudaStream); + recHits.flags = cms::cuda::make_device_unique(size, cudaStream); + } + }; + } //namespace weights +} //namespace ecal + +#endif // RecoLocalCalo_EcalRecProducers_plugins_DeclsForKernelsPhase2_h From 67d34665beeb7218d5076e89f9d77aa2dfc4e664 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:13:34 +0100 Subject: [PATCH 12/23] Changed converter to use switch description --- .../plugins/EcalCPUUncalibRecHitProducer.cc | 38 +++++++++---------- 1 file changed, 17 insertions(+), 21 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc index cc29e24e8e9e9..509eea5a3ddb5 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc @@ -9,6 +9,7 @@ #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/MakerMacros.h" +#include "FWCore/ParameterSet/interface/EmptyGroupDescription.h" // algorithm specific @@ -27,14 +28,14 @@ class EcalCPUUncalibRecHitProducer : public edm::stream::EDProducer>; - edm::EDGetTokenT recHitsInEBToken_, recHitsInEEToken_; + const edm::EDGetTokenT recHitsInEBToken_, recHitsInEEToken_; using OutputProduct = ecal::UncalibratedRecHit>; - edm::EDPutTokenT recHitsOutEBToken_, recHitsOutEEToken_; + const edm::EDPutTokenT recHitsOutEBToken_, recHitsOutEEToken_; OutputProduct recHitsEB_, recHitsEE_; - bool containsTimingInformation_; + const bool containsTimingInformation_; }; void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptions& confDesc) { @@ -42,12 +43,13 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio desc.add("recHitsInLabelEB", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEB"}); desc.add("recHitsOutLabelEB", "EcalUncalibRecHitsEB"); - desc.add("containsTimingInformation", false); - desc.add("isPhase2", false); - - desc.add("recHitsInLabelEE", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"}); - desc.add("recHitsOutLabelEE", "EcalUncalibRecHitsEE"); + desc.ifValue( + edm::ParameterDescription("isPhase2", false, true), + false >> (edm::ParameterDescription( + "recHitsInLabelEE", edm::InputTag{"ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"}, true) and + edm::ParameterDescription("recHitsOutLabelEE", "EcalUncalibRecHitsEE", true)) or + true >> edm::EmptyGroupDescription()); confDesc.add("ecalCPUUncalibRecHitProducer", desc); } @@ -55,13 +57,12 @@ void EcalCPUUncalibRecHitProducer::fillDescriptions(edm::ConfigurationDescriptio EcalCPUUncalibRecHitProducer::EcalCPUUncalibRecHitProducer(const edm::ParameterSet& ps) : isPhase2_{ps.getParameter("isPhase2")}, recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, + recHitsInEEToken_{isPhase2_ ? edm::EDGetTokenT{} + : consumes(ps.getParameter("recHitsInLabelEE"))}, recHitsOutEBToken_{produces(ps.getParameter("recHitsOutLabelEB"))}, - containsTimingInformation_{ps.getParameter("containsTimingInformation")} { - if (!isPhase2_) { - recHitsInEEToken_ = consumes(ps.getParameter("recHitsInLabelEE")); - recHitsOutEEToken_ = produces(ps.getParameter("recHitsOutLabelEE")); - } -} + recHitsOutEEToken_{isPhase2_ ? edm::EDPutTokenT{} + : produces(ps.getParameter("recHitsOutLabelEE"))}, + containsTimingInformation_{ps.getParameter("containsTimingInformation")} {} EcalCPUUncalibRecHitProducer::~EcalCPUUncalibRecHitProducer() {} @@ -84,11 +85,6 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream())); }; - if (!isPhase2_) { - auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); - auto const& eeRecHits = ctx.get(eeRecHitsProduct); - recHitsEE_.resize(eeRecHits.size); - } // enqeue transfers lambdaToTransfer(recHitsEB_.did, ebRecHits.did.get()); lambdaToTransfer(recHitsEB_.amplitudesAll, ebRecHits.amplitudesAll.get()); @@ -132,4 +128,4 @@ void EcalCPUUncalibRecHitProducer::produce(edm::Event& event, edm::EventSetup co } } -DEFINE_FWK_MODULE(EcalCPUUncalibRecHitProducer); \ No newline at end of file +DEFINE_FWK_MODULE(EcalCPUUncalibRecHitProducer); From f9edddeb6d405d12331e080a3eaaa7f9cd2f21a6 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:14:49 +0100 Subject: [PATCH 13/23] Added Phase2 CPU to GPU converter --- .../plugins/EcalPhase2DigiToGPUProducer.cc | 100 ++++++++++++++++++ 1 file changed, 100 insertions(+) create mode 100644 RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc new file mode 100644 index 0000000000000..f5328eca17701 --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc @@ -0,0 +1,100 @@ +#include "FWCore/Framework/interface/stream/EDProducer.h" +#include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/EventSetup.h" +#include "FWCore/Framework/interface/MakerMacros.h" + +#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" +#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" +#include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" + +#include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" +#include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" + +#include "DeclsForKernelsPhase2.h" + +class EcalPhase2DigiToGPUProducer : public edm::stream::EDProducer { +public: + explicit EcalPhase2DigiToGPUProducer(const edm::ParameterSet& ps); + ~EcalPhase2DigiToGPUProducer() override = default; + static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + + void acquire(edm::Event const& event, edm::EventSetup const& setup, edm::WaitingTaskWithArenaHolder holder) override; + void produce(edm::Event& evt, edm::EventSetup const& setup) override; + +private: + const edm::EDGetTokenT digiCollectionToken_; + const edm::EDPutTokenT>> + digisCollectionToken_; + uint32_t size_; + + ecal::DigisCollection<::calo::common::DevStoragePolicy> digis_; + + cms::cuda::ContextState cudaState_; +}; + +void EcalPhase2DigiToGPUProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { + edm::ParameterSetDescription desc; + + desc.add("BarrelDigis", edm::InputTag("simEcalUnsuppressedDigis", "")); + desc.add("digisLabelEB", "ebDigis"); + + descriptions.addWithDefaultLabel(desc); +} + +EcalPhase2DigiToGPUProducer::EcalPhase2DigiToGPUProducer(const edm::ParameterSet& ps) + : digiCollectionToken_(consumes(ps.getParameter("BarrelDigis"))), + digisCollectionToken_(produces>>( + ps.getParameter("digisLabelEB"))) {} + +void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, + edm::EventSetup const& setup, + edm::WaitingTaskWithArenaHolder holder) { + cms::cuda::ScopedContextAcquire ctx{event.streamID(), std::move(holder), cudaState_}; + + //input data from event + const auto& pdigis = event.get(digiCollectionToken_); + + size_ = pdigis.size(); + + //allocate device pointers for output + digis_.ids = cms::cuda::make_device_unique(size_, ctx.stream()); + digis_.data = cms::cuda::make_device_unique(size_ * EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); + + //allocate host pointers for holding product data and id vectors + auto idstmp = cms::cuda::make_host_unique(size_,ctx.stream()); + auto datatmp = cms::cuda::make_host_unique(size_ * EcalDataFrame_Ph2::MAXSAMPLES,ctx.stream()); + + //iterate over digis + uint32_t i = 0; + for (const auto& pdigi : pdigis) { + const int nSamples = pdigi.size(); + //assign id to output vector + idstmp.get()[i] = pdigi.id(); + //iterate over sample in digi + for (int sample = 0; sample < nSamples; ++sample) { + //get samples from input digi + EcalLiteDTUSample thisSample = pdigi[sample]; + //assign adc data to output + datatmp.get()[i * nSamples + sample] = thisSample.raw(); + } + ++i; + } + + //copy output vectors into member variable device pointers for the output struct + + cudaCheck(cudaMemcpyAsync( + digis_.ids.get(), idstmp.get(), size_ * sizeof(uint32_t), cudaMemcpyHostToDevice, ctx.stream())); + cudaCheck(cudaMemcpyAsync( + digis_.data.get(), datatmp.get(), size_ * EcalDataFrame_Ph2::MAXSAMPLES * sizeof(uint16_t), cudaMemcpyHostToDevice, ctx.stream())); +} + +void EcalPhase2DigiToGPUProducer::produce(edm::Event& event, edm::EventSetup const& setup) { + //get cuda context state for producer + cms::cuda::ScopedContextProduce ctx{cudaState_}; + digis_.size = size_; + + //emplace output in the context + ctx.emplace(event, digisCollectionToken_, std::move(digis_)); +} + +DEFINE_FWK_MODULE(EcalPhase2DigiToGPUProducer); From b59b8c1cd801324921b93a75a3649706495ee229 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:15:52 +0100 Subject: [PATCH 14/23] Changed converter to use switch description --- .../EcalUncalibRecHitConvertGPU2CPUFormat.cc | 33 ++++++++++--------- 1 file changed, 17 insertions(+), 16 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc index 941da16f26717..4ddfd56346dee 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc @@ -7,6 +7,7 @@ #include "FWCore/Framework/interface/MakerMacros.h" #include "FWCore/Framework/interface/stream/EDProducer.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/ParameterSet/interface/EmptyGroupDescription.h" class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { public: @@ -19,39 +20,39 @@ class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { void produce(edm::Event&, edm::EventSetup const&) override; private: - bool isPhase2_; + const bool isPhase2_; const edm::EDGetTokenT recHitsGPUEB_; - edm::EDGetTokenT recHitsGPUEE_; + const edm::EDGetTokenT recHitsGPUEE_; const std::string recHitsLabelCPUEB_; - std::string recHitsLabelCPUEE_; + const std::string recHitsLabelCPUEE_; }; void EcalUncalibRecHitConvertGPU2CPUFormat::fillDescriptions(edm::ConfigurationDescriptions& confDesc) { edm::ParameterSetDescription desc; desc.add("recHitsLabelGPUEB", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEB")); - desc.add("recHitsLabelCPUEB", "EcalUncalibRecHitsEB"); - - desc.add("isPhase2", false); - - desc.add("recHitsLabelGPUEE", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE")); - desc.add("recHitsLabelCPUEE", "EcalUncalibRecHitsEE"); - + desc.ifValue( + edm::ParameterDescription("isPhase2", false, true), + false >> + (edm::ParameterDescription( + "recHitsLabelGPUEE", edm::InputTag("ecalUncalibRecHitProducerGPU", "EcalUncalibRecHitsEE"), true) and + edm::ParameterDescription("recHitsLabelCPUEE", "EcalUncalibRecHitsEE", true)) or + true >> edm::EmptyGroupDescription()); confDesc.add("ecalUncalibRecHitConvertGPU2CPUFormat", desc); } EcalUncalibRecHitConvertGPU2CPUFormat::EcalUncalibRecHitConvertGPU2CPUFormat(const edm::ParameterSet& ps) : isPhase2_{ps.getParameter("isPhase2")}, recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, - recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")} { + recHitsGPUEE_{isPhase2_ ? edm::EDGetTokenT{} + : consumes(ps.getParameter("recHitsLabelGPUEE"))}, + recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")}, + recHitsLabelCPUEE_{isPhase2_ ? std::string{""} : ps.getParameter("recHitsLabelCPUEE")} { produces(recHitsLabelCPUEB_); - if (!isPhase2_) { - recHitsGPUEE_ = consumes(ps.getParameter("recHitsLabelGPUEE")); - recHitsLabelCPUEE_ = ps.getParameter("recHitsLabelCPUEE"); + if (!isPhase2_) produces(recHitsLabelCPUEE_); - } } EcalUncalibRecHitConvertGPU2CPUFormat::~EcalUncalibRecHitConvertGPU2CPUFormat() {} @@ -97,4 +98,4 @@ void EcalUncalibRecHitConvertGPU2CPUFormat::produce(edm::Event& event, edm::Even event.put(std::move(recHitsCPUEB), recHitsLabelCPUEB_); } -DEFINE_FWK_MODULE(EcalUncalibRecHitConvertGPU2CPUFormat); \ No newline at end of file +DEFINE_FWK_MODULE(EcalUncalibRecHitConvertGPU2CPUFormat); From 487051f7c70474c9858f6eadc46eb65b1fdac42c Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:16:55 +0100 Subject: [PATCH 15/23] Code checks --- .../plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu index 20e9849b3959f..e33daf81ed362 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu @@ -32,4 +32,4 @@ namespace ecal { } } // namespace weights -} // namespace ecal \ No newline at end of file +} // namespace ecal From c51e5791eb2970b4f0919d371c4bfa8944c12d47 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:17:46 +0100 Subject: [PATCH 16/23] Changed Ph2 to Phase2 --- .../plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h index 5519879c58479..d3baf63345b4e 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h @@ -3,7 +3,7 @@ #include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" -#include "DeclsForKernelsPh2.h" +#include "DeclsForKernelsPhase2.h" namespace ecal { namespace weights { @@ -16,4 +16,4 @@ namespace ecal { } // namespace weights } // namespace ecal -#endif // RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsAlgoGPU_h \ No newline at end of file +#endif // RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsAlgoGPU_h From fcd98f86f446dbab7eed5b59deeba16addfe0227 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:18:43 +0100 Subject: [PATCH 17/23] Changed Ph2 to Phase2 --- .../plugins/EcalUncalibRecHitPhase2WeightsKernels.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu index 04adbfb3452cf..9cbdb879a4a3e 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu @@ -28,7 +28,7 @@ namespace ecal { extern __shared__ char shared_mem[]; double* shr_weights = (double*)&shared_mem[0]; float* shr_amp = (float*)&shared_mem[nsamples * sizeof(double)]; - uint16_t* shr_digis = (uint16_t*)&shared_mem[nsamples * sizeof(double)+ nchannels_per_block * sizeof(float)]; + uint16_t* shr_digis = (uint16_t*)&shared_mem[nsamples * sizeof(double) + nchannels_per_block * sizeof(float)]; shr_weights = weights; @@ -45,8 +45,8 @@ namespace ecal { for (int sample = 0; sample < nsamples; ++sample) { const unsigned int idx = threadIdx.x * nsamples + sample; const auto shr_digi = shr_digis[idx]; - shr_amp[btx] += (static_cast(ecalLiteDTU::adc(shr_digi)) * ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * - shr_weights[sample]); + shr_amp[btx] += (static_cast(ecalLiteDTU::adc(shr_digi)) * + ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * shr_weights[sample]); } amplitude[tx] = shr_amp[btx]; amplitudeError[tx] = 1.0f; @@ -59,4 +59,4 @@ namespace ecal { } //if within nchannels } //kernel } //namespace weights -} //namespace ecal \ No newline at end of file +} //namespace ecal From 1dcac8f8e622f6b68a8fa7ae4785a8f07c9dea52 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:19:24 +0100 Subject: [PATCH 18/23] Changed Ph2 to Phase2 --- .../plugins/EcalUncalibRecHitPhase2WeightsKernels.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h index 4517b75d53d1f..f5d26bd5364dc 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h @@ -1,7 +1,7 @@ #ifndef RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsKernels_h #define RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsKernels_h -#include "DeclsForKernelsPh2.h" +#include "DeclsForKernelsPhase2.h" namespace ecal { namespace weights { @@ -17,4 +17,4 @@ namespace ecal { } //namespace weights } //namespace ecal -#endif \ No newline at end of file +#endif From 722ef6903fcb19dd95c33761cedb7647026f94b2 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:20:40 +0100 Subject: [PATCH 19/23] Changed weights_ to use cms::cuda::HostAllocator --- ...alUncalibRecHitPhase2WeightsProducerGPU.cc | 25 ++++++++++--------- 1 file changed, 13 insertions(+), 12 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc index 254bc94c3f619..ec5a031fd24b4 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc @@ -4,11 +4,12 @@ #include "FWCore/Framework/interface/MakerMacros.h" #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" +#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" #include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" #include "EcalUncalibRecHitPhase2WeightsAlgoGPU.h" -#include "DeclsForKernelsPh2.h" +#include "DeclsForKernelsPhase2.h" class EcalUncalibRecHitPhase2WeightsProducerGPU : public edm::stream::EDProducer { public: @@ -21,7 +22,7 @@ class EcalUncalibRecHitPhase2WeightsProducerGPU : public edm::stream::EDProducer void produce(edm::Event &, edm::EventSetup const &) override; private: - const std::vector weights_; + const std::vector> weights_; using InputProduct = cms::cuda::Product>; const edm::EDGetTokenT digisToken_; @@ -33,12 +34,15 @@ class EcalUncalibRecHitPhase2WeightsProducerGPU : public edm::stream::EDProducer cms::cuda::ContextState cudaState_; - uint32_t n_; + uint32_t size_; }; // constructor with initialisation of elements EcalUncalibRecHitPhase2WeightsProducerGPU::EcalUncalibRecHitPhase2WeightsProducerGPU(const edm::ParameterSet &ps) - : weights_(ps.getParameter>("weights")), + : // use lambda to initialise the vector with CUDA::HostAllocator from a normal vector + weights_([tmp = ps.getParameter>("weights")] { + return std::vector>(tmp.begin(), tmp.end()); + }()), digisToken_{consumes(ps.getParameter("digisLabelEB"))}, recHitsToken_{produces(ps.getParameter("recHitsLabelEB"))} {} @@ -81,16 +85,13 @@ void EcalUncalibRecHitPhase2WeightsProducerGPU::acquire(edm::Event const &event, // get actual obj auto const &digis = ctx.get(digisProduct); - n_ = digis.size; + size_ = digis.size; // if no digis stop here - if (n_ == 0) + if (size_ == 0) return; - // weights to GPU - - cms::cuda::device::unique_ptr weights_d = - cms::cuda::make_device_unique(EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); + auto weights_d = cms::cuda::make_device_unique(EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); cudaCheck(cudaMemcpyAsync(weights_d.get(), weights_.data(), @@ -99,7 +100,7 @@ void EcalUncalibRecHitPhase2WeightsProducerGPU::acquire(edm::Event const &event, ctx.stream())); // output on GPU - eventOutputDataGPU_.allocate(n_, ctx.stream()); + eventOutputDataGPU_.allocate(size_, ctx.stream()); ecal::weights::entryPoint(digis, eventOutputDataGPU_, weights_d, ctx.stream()); } @@ -108,7 +109,7 @@ void EcalUncalibRecHitPhase2WeightsProducerGPU::produce(edm::Event &event, const cms::cuda::ScopedContextProduce ctx{cudaState_}; // set the size of digis - eventOutputDataGPU_.recHits.size = n_; + eventOutputDataGPU_.recHits.size = size_; // put into the event ctx.emplace(event, recHitsToken_, std::move(eventOutputDataGPU_.recHits)); From b1ba0842a97ee0add91c7130315ef94a7bb9292b Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever <89704887+ChrisSandever@users.noreply.github.com> Date: Mon, 30 May 2022 23:22:15 +0100 Subject: [PATCH 20/23] Chanhed Ph2 to Phase2 and removes EE objects --- .../python/ecalUncalibRecHitPhase2_cff.py | 27 ++++++++++--------- 1 file changed, 15 insertions(+), 12 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py index 7b03dbf9bee92..5b5f37d8905ea 100644 --- a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py +++ b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py @@ -14,39 +14,42 @@ # conditions used on gpu -from RecoLocalCalo.EcalRecProducers.ecalPh2DigiToGPUProducer_cfi import ecalPh2DigiToGPUProducer as _ecalPh2DigiToGPUProducer -ecalPh2DigiToGPUProducer = _ecalPh2DigiToGPUProducer.clone() +from RecoLocalCalo.EcalRecProducers.ecalPhase2DigiToGPUProducer_cfi import ecalPhase2DigiToGPUProducer as _ecalPhase2DigiToGPUProducer +ecalPhase2DigiToGPUProducer = _ecalPhase2DigiToGPUProducer.clone() # gpu weights from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2GPU_cfi import ecalUncalibRecHitPhase2GPU as _ecalUncalibRecHitPhase2GPU ecalUncalibRecHitPhase2GPU = _ecalUncalibRecHitPhase2GPU.clone( - digisLabelEB = cms.InputTag('ecalPh2DigiToGPUProducer', 'ebDigis') + digisLabelEB = ('ecalPhase2DigiToGPUProducer', 'ebDigis') ) # copy the uncalibrated rechits from GPU to CPU from RecoLocalCalo.EcalRecProducers.ecalCPUUncalibRecHitProducer_cfi import ecalCPUUncalibRecHitProducer as _ecalCPUUncalibRecHitProducer -ecalMultiFitUncalibRecHitSoAnew = _ecalCPUUncalibRecHitProducer.clone( - recHitsInLabelEB = cms.InputTag('ecalUncalibRecHitPhase2GPU', 'EcalUncalibRecHitsEB'), - isPhase2 = cms.bool(True) +ecalUncalibRecHitSoA = _ecalCPUUncalibRecHitProducer.clone( + recHitsInLabelEB = ('ecalUncalibRecHitPhase2GPU', 'EcalUncalibRecHitsEB'), + isPhase2 = True, + recHitsInLabelEE = None, # remove unneeded Phase1 parameters + recHitsOutLabelEE = None ) -# convert the uncalibrated rechits from SoA to legacy format from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitConvertGPU2CPUFormat_cfi import ecalUncalibRecHitConvertGPU2CPUFormat as _ecalUncalibRecHitConvertGPU2CPUFormat gpu.toModify(ecalUncalibRecHitPhase2, - cuda = _ecalUncalibRecHitConvertGPU2CPUFormat.clone( - recHitsLabelGPUEB = cms.InputTag('ecalMultiFitUncalibRecHitSoAnew', 'EcalUncalibRecHitsEB'), - isPhase2 = cms.bool(True) + cuda = _ecalUncalibRecHitConvertGPU2CPUFormat.clone( + isPhase2 = cms.bool(True), + recHitsLabelGPUEB = cms.InputTag('ecalUncalibRecHitSoA', 'EcalUncalibRecHitsEB'), + recHitsLabelGPUEE = None, # remove unneeded Phase1 parameters + recHitsLabelCPUEE = None ) ) gpu.toReplaceWith(ecalUncalibRecHitPhase2Task, cms.Task( # convert phase2 digis to GPU SoA - ecalPh2DigiToGPUProducer, + ecalPhase2DigiToGPUProducer, # ECAL weights running on GPU ecalUncalibRecHitPhase2GPU, # copy the uncalibrated rechits from GPU to CPU - ecalMultiFitUncalibRecHitSoAnew, + ecalUncalibRecHitSoA, # ECAL multifit running on CPU, or convert the uncalibrated rechits from SoA to legacy format ecalUncalibRecHitPhase2, )) From e48b02d1bd2aaacccde6ed6b99bc3a8f3313a668 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever Date: Tue, 31 May 2022 01:07:42 +0200 Subject: [PATCH 21/23] Code checks --- .../plugins/EcalPhase2DigiToGPUProducer.cc | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc index f5328eca17701..7d29cf16d5d68 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc @@ -47,8 +47,8 @@ EcalPhase2DigiToGPUProducer::EcalPhase2DigiToGPUProducer(const edm::ParameterSet ps.getParameter("digisLabelEB"))) {} void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, - edm::EventSetup const& setup, - edm::WaitingTaskWithArenaHolder holder) { + edm::EventSetup const& setup, + edm::WaitingTaskWithArenaHolder holder) { cms::cuda::ScopedContextAcquire ctx{event.streamID(), std::move(holder), cudaState_}; //input data from event @@ -61,8 +61,8 @@ void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, digis_.data = cms::cuda::make_device_unique(size_ * EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); //allocate host pointers for holding product data and id vectors - auto idstmp = cms::cuda::make_host_unique(size_,ctx.stream()); - auto datatmp = cms::cuda::make_host_unique(size_ * EcalDataFrame_Ph2::MAXSAMPLES,ctx.stream()); + auto idstmp = cms::cuda::make_host_unique(size_, ctx.stream()); + auto datatmp = cms::cuda::make_host_unique(size_ * EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); //iterate over digis uint32_t i = 0; @@ -82,10 +82,13 @@ void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, //copy output vectors into member variable device pointers for the output struct - cudaCheck(cudaMemcpyAsync( - digis_.ids.get(), idstmp.get(), size_ * sizeof(uint32_t), cudaMemcpyHostToDevice, ctx.stream())); - cudaCheck(cudaMemcpyAsync( - digis_.data.get(), datatmp.get(), size_ * EcalDataFrame_Ph2::MAXSAMPLES * sizeof(uint16_t), cudaMemcpyHostToDevice, ctx.stream())); + cudaCheck( + cudaMemcpyAsync(digis_.ids.get(), idstmp.get(), size_ * sizeof(uint32_t), cudaMemcpyHostToDevice, ctx.stream())); + cudaCheck(cudaMemcpyAsync(digis_.data.get(), + datatmp.get(), + size_ * EcalDataFrame_Ph2::MAXSAMPLES * sizeof(uint16_t), + cudaMemcpyHostToDevice, + ctx.stream())); } void EcalPhase2DigiToGPUProducer::produce(edm::Event& event, edm::EventSetup const& setup) { From be0268c63905d3768c9e614541e13fe5dc409926 Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever Date: Tue, 31 May 2022 11:15:53 +0200 Subject: [PATCH 22/23] Made else statement in converter --- .../EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc index 509eea5a3ddb5..60600b0376d3e 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc @@ -96,10 +96,10 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, lambdaToTransfer(recHitsEB_.jitter, ebRecHits.jitter.get()); lambdaToTransfer(recHitsEB_.jitterError, ebRecHits.jitterError.get()); } - if (isPhase2_) + if (isPhase2_) { lambdaToTransfer(recHitsEB_.amplitudeError, ebRecHits.amplitudeError.get()); - if (!isPhase2_) { + } else { auto const& eeRecHitsProduct = event.get(recHitsInEEToken_); auto const& eeRecHits = ctx.get(eeRecHitsProduct); recHitsEE_.resize(eeRecHits.size); From 0820f197b5d49d5b94f26f1efcbe9199255a7c4b Mon Sep 17 00:00:00 2001 From: Christopher Rhys Sandever Date: Tue, 5 Jul 2022 20:27:00 +0200 Subject: [PATCH 23/23] Reverted kernel changes and added comments --- .../plugins/EcalPhase2DigiToGPUProducer.cc | 2 +- .../EcalUncalibRecHitPhase2WeightsAlgoGPU.cu | 18 ++++----- .../EcalUncalibRecHitPhase2WeightsAlgoGPU.h | 8 ++-- .../EcalUncalibRecHitPhase2WeightsKernels.cu | 38 +++++++++---------- .../EcalUncalibRecHitPhase2WeightsKernels.h | 2 +- .../EcalUncalibRecHitPhase2WeightsProducer.cc | 1 + ...alUncalibRecHitPhase2WeightsProducerGPU.cc | 3 +- 7 files changed, 36 insertions(+), 36 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc index 7d29cf16d5d68..fc0b6c8a41dfe 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc @@ -56,6 +56,7 @@ void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, size_ = pdigis.size(); + digis_.size = size_; //allocate device pointers for output digis_.ids = cms::cuda::make_device_unique(size_, ctx.stream()); digis_.data = cms::cuda::make_device_unique(size_ * EcalDataFrame_Ph2::MAXSAMPLES, ctx.stream()); @@ -94,7 +95,6 @@ void EcalPhase2DigiToGPUProducer::acquire(edm::Event const& event, void EcalPhase2DigiToGPUProducer::produce(edm::Event& event, edm::EventSetup const& setup) { //get cuda context state for producer cms::cuda::ScopedContextProduce ctx{cudaState_}; - digis_.size = size_; //emplace output in the context ctx.emplace(event, digisCollectionToken_, std::move(digis_)); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu index e33daf81ed362..0ed3ede25734e 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu @@ -7,18 +7,16 @@ namespace ecal { namespace weights { - void entryPoint(ecal::DigisCollection const& digis, - EventOutputDataGPU& eventOutputGPU, - cms::cuda::device::unique_ptr& weights_d, - cudaStream_t cudaStream) { - unsigned int totalChannels = digis.size; + void phase2Weights(ecal::DigisCollection const& digis, + EventOutputDataGPU& eventOutputGPU, + cms::cuda::device::unique_ptr& weights_d, + cudaStream_t cudaStream) { + unsigned int const totalChannels = digis.size; // 64 threads per block best occupancy from Nsight compute profiler - unsigned int nchannels_per_block = 64; - unsigned int threads_1d = nchannels_per_block; - unsigned int blocks_1d = (totalChannels / threads_1d) + 1; - // shared bytes from size of weight constants, digi samples per block, uncalib rechits amplitudes per block + unsigned int const threads_1d = 64; + unsigned int const blocks_1d = (totalChannels + threads_1d - 1) / threads_1d; int shared_bytes = EcalDataFrame_Ph2::MAXSAMPLES * sizeof(double) + - nchannels_per_block * (EcalDataFrame_Ph2::MAXSAMPLES * (sizeof(uint16_t)) + sizeof(float)); + threads_1d * (EcalDataFrame_Ph2::MAXSAMPLES * (sizeof(uint16_t)) + sizeof(float)); Phase2WeightsKernel<<>>( digis.data.get(), digis.ids.get(), diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h index d3baf63345b4e..82e0ebca59f32 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h @@ -8,10 +8,10 @@ namespace ecal { namespace weights { - void entryPoint(ecal::DigisCollection const&, - EventOutputDataGPU&, - cms::cuda::device::unique_ptr&, - cudaStream_t); + void phase2Weights(ecal::DigisCollection const&, + EventOutputDataGPU&, + cms::cuda::device::unique_ptr&, + cudaStream_t); } // namespace weights } // namespace ecal diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu index 9cbdb879a4a3e..f2fcad4510d82 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu @@ -11,51 +11,51 @@ namespace ecal { namespace weights { __global__ void Phase2WeightsKernel(uint16_t const* digis_in, - uint32_t const* dids, - ::ecal::reco::StorageScalarType* amplitude, - ::ecal::reco::StorageScalarType* amplitudeError, - uint32_t* dids_out, + uint32_t const* __restrict__ dids, + ::ecal::reco::StorageScalarType* __restrict__ amplitude, + ::ecal::reco::StorageScalarType* __restrict__ amplitudeError, + uint32_t* __restrict__ dids_out, int const nchannels, - double* weights, - uint32_t* flags) { + double const* __restrict__ weights, + uint32_t* __restrict__ flags) { constexpr int nsamples = EcalDataFrame_Ph2::MAXSAMPLES; int const tx = threadIdx.x + blockIdx.x * blockDim.x; unsigned int nchannels_per_block = blockDim.x; + unsigned int const threadx = threadIdx.x; if (tx < nchannels) { - auto const did = DetId{dids[tx]}; - //dynamic shared memory extern __shared__ char shared_mem[]; double* shr_weights = (double*)&shared_mem[0]; float* shr_amp = (float*)&shared_mem[nsamples * sizeof(double)]; uint16_t* shr_digis = (uint16_t*)&shared_mem[nsamples * sizeof(double) + nchannels_per_block * sizeof(float)]; + for (int i = 0; i < nsamples; ++i) + shr_weights[i] = weights[i]; - shr_weights = weights; - - unsigned int bx = blockIdx.x; //block index - unsigned int btx = threadIdx.x; + unsigned int const bx = blockIdx.x; //block index for (int sample = 0; sample < nsamples; ++sample) { - const unsigned int idx = threadIdx.x * nsamples + sample; + int const idx = threadx * nsamples + sample; shr_digis[idx] = digis_in[bx * nchannels_per_block * nsamples + idx]; } + shr_amp[threadx] = 0.0; + __syncthreads(); - shr_amp[btx] = 0.0; + auto const did = DetId{dids[tx]}; CMS_UNROLL_LOOP for (int sample = 0; sample < nsamples; ++sample) { const unsigned int idx = threadIdx.x * nsamples + sample; const auto shr_digi = shr_digis[idx]; - shr_amp[btx] += (static_cast(ecalLiteDTU::adc(shr_digi)) * - ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * shr_weights[sample]); + shr_amp[threadx] += (static_cast(ecalLiteDTU::adc(shr_digi)) * + ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * shr_weights[sample]); } - amplitude[tx] = shr_amp[btx]; + const unsigned int tdx = threadIdx.x * nsamples; + amplitude[tx] = shr_amp[threadx]; amplitudeError[tx] = 1.0f; dids_out[tx] = did.rawId(); flags[tx] = 0; - if (ecalLiteDTU::gainId(shr_digis[btx * nsamples + nsamples - 1])) { + if (ecalLiteDTU::gainId(shr_digis[tdx + nsamples - 1])) { flags[tx] = 0x1 << EcalUncalibratedRecHit::kHasSwitchToGain1; } - } //if within nchannels } //kernel } //namespace weights diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h index f5d26bd5364dc..22fd89eccdefa 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h @@ -12,7 +12,7 @@ namespace ecal { ::ecal::reco::StorageScalarType* amplitudeErrorEB, uint32_t* dids_outEB, int const nchannels, - double* weights_d, + double const* weights_d, uint32_t* flagsEB); } //namespace weights } //namespace ecal diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducer.cc index 414ad351c0236..ae67e128d5644 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducer.cc @@ -33,6 +33,7 @@ void EcalUncalibRecHitPhase2WeightsProducer::fillDescriptions(edm::Configuration desc.add("EBhitCollection", "EcalUncalibRecHitsEB"); desc.add("tRise", 0.2); desc.add("tFall", 2.); + // The below weights values should be kept up to date with GPU version of this module desc.add>("weights", {-0.121016, -0.119899, diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc index ec5a031fd24b4..f324b71799e79 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc @@ -50,6 +50,7 @@ void EcalUncalibRecHitPhase2WeightsProducerGPU::fillDescriptions(edm::Configurat edm::ParameterSetDescription desc; desc.add("recHitsLabelEB", "EcalUncalibRecHitsEB"); + //The below weights values should be kept up to date with those on the CPU version of this module desc.add>("weights", {-0.121016, -0.119899, @@ -102,7 +103,7 @@ void EcalUncalibRecHitPhase2WeightsProducerGPU::acquire(edm::Event const &event, // output on GPU eventOutputDataGPU_.allocate(size_, ctx.stream()); - ecal::weights::entryPoint(digis, eventOutputDataGPU_, weights_d, ctx.stream()); + ecal::weights::phase2Weights(digis, eventOutputDataGPU_, weights_d, ctx.stream()); } void EcalUncalibRecHitPhase2WeightsProducerGPU::produce(edm::Event &event, const edm::EventSetup &setup) {