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); 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 diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalCPUUncalibRecHitProducer.cc index 801d378c7c391..60600b0376d3e 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,32 +28,40 @@ 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) { 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.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); } EcalCPUUncalibRecHitProducer::EcalCPUUncalibRecHitProducer(const edm::ParameterSet& ps) - : recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, - recHitsInEEToken_{consumes(ps.getParameter("recHitsInLabelEE"))}, + : isPhase2_{ps.getParameter("isPhase2")}, + recHitsInEBToken_{consumes(ps.getParameter("recHitsInLabelEB"))}, + recHitsInEEToken_{isPhase2_ ? edm::EDGetTokenT{} + : consumes(ps.getParameter("recHitsInLabelEE"))}, recHitsOutEBToken_{produces(ps.getParameter("recHitsOutLabelEB"))}, - recHitsOutEEToken_{produces(ps.getParameter("recHitsOutLabelEE"))}, + recHitsOutEEToken_{isPhase2_ ? edm::EDPutTokenT{} + : produces(ps.getParameter("recHitsOutLabelEE"))}, containsTimingInformation_{ps.getParameter("containsTimingInformation")} {} EcalCPUUncalibRecHitProducer::~EcalCPUUncalibRecHitProducer() {} @@ -62,14 +71,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; @@ -81,40 +87,45 @@ void EcalCPUUncalibRecHitProducer::acquire(edm::Event const& event, // 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_.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 (isPhase2_) { + lambdaToTransfer(recHitsEB_.amplitudeError, ebRecHits.amplitudeError.get()); + + } else { + 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_.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 (!isPhase2_) { + auto recHitsOutEE = std::make_unique(std::move(recHitsEE_)); + event.put(recHitsOutEEToken_, std::move(recHitsOutEE)); + } } DEFINE_FWK_MODULE(EcalCPUUncalibRecHitProducer); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc new file mode 100644 index 0000000000000..fc0b6c8a41dfe --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalPhase2DigiToGPUProducer.cc @@ -0,0 +1,103 @@ +#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(); + + 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()); + + //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_}; + + //emplace output in the context + ctx.emplace(event, digisCollectionToken_, std::move(digis_)); +} + +DEFINE_FWK_MODULE(EcalPhase2DigiToGPUProducer); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc index b26fbe3a0c572..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,73 +20,82 @@ class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { void produce(edm::Event&, edm::EventSetup const&) override; private: + const bool isPhase2_; const edm::EDGetTokenT recHitsGPUEB_; const edm::EDGetTokenT recHitsGPUEE_; - const std::string recHitsLabelCPUEB_, recHitsLabelCPUEE_; + const std::string recHitsLabelCPUEB_; + const 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("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) - : recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, - recHitsGPUEE_{consumes(ps.getParameter("recHitsLabelGPUEE"))}, + : isPhase2_{ps.getParameter("isPhase2")}, + recHitsGPUEB_{consumes(ps.getParameter("recHitsLabelGPUEB"))}, + recHitsGPUEE_{isPhase2_ ? edm::EDGetTokenT{} + : consumes(ps.getParameter("recHitsLabelGPUEE"))}, recHitsLabelCPUEB_{ps.getParameter("recHitsLabelCPUEB")}, - recHitsLabelCPUEE_{ps.getParameter("recHitsLabelCPUEE")} { + recHitsLabelCPUEE_{isPhase2_ ? std::string{""} : ps.getParameter("recHitsLabelCPUEE")} { produces(recHitsLabelCPUEB_); - produces(recHitsLabelCPUEE_); + if (!isPhase2_) + 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]); + 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, 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 (!isPhase2_) { + 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]); + (*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); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu new file mode 100644 index 0000000000000..0ed3ede25734e --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu @@ -0,0 +1,33 @@ +#include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" +#include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" + +#include "EcalUncalibRecHitPhase2WeightsKernels.h" +#include "EcalUncalibRecHitPhase2WeightsAlgoGPU.h" + +namespace ecal { + namespace weights { + + 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 const threads_1d = 64; + unsigned int const blocks_1d = (totalChannels + threads_1d - 1) / threads_1d; + int shared_bytes = EcalDataFrame_Ph2::MAXSAMPLES * sizeof(double) + + threads_1d * (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 diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.h new file mode 100644 index 0000000000000..82e0ebca59f32 --- /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 "DeclsForKernelsPhase2.h" + +namespace ecal { + namespace weights { + + void phase2Weights(ecal::DigisCollection const&, + EventOutputDataGPU&, + cms::cuda::device::unique_ptr&, + cudaStream_t); + + } // namespace weights +} // namespace ecal + +#endif // RecoLocalCalo_EcalRecProducers_plugins_EcalUncalibRecHitPhase2WeightsAlgoGPU_h diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu new file mode 100644 index 0000000000000..f2fcad4510d82 --- /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* __restrict__ dids, + ::ecal::reco::StorageScalarType* __restrict__ amplitude, + ::ecal::reco::StorageScalarType* __restrict__ amplitudeError, + uint32_t* __restrict__ dids_out, + int const nchannels, + 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) { + 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]; + + unsigned int const bx = blockIdx.x; //block index + + for (int sample = 0; sample < 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(); + + 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[threadx] += (static_cast(ecalLiteDTU::adc(shr_digi)) * + ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * shr_weights[sample]); + } + 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[tdx + nsamples - 1])) { + flags[tx] = 0x1 << EcalUncalibratedRecHit::kHasSwitchToGain1; + } + } //if within nchannels + } //kernel + } //namespace weights +} //namespace ecal diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.h new file mode 100644 index 0000000000000..22fd89eccdefa --- /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 "DeclsForKernelsPhase2.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 const* weights_d, + uint32_t* flagsEB); + } //namespace weights +} //namespace ecal + +#endif 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 new file mode 100644 index 0000000000000..f324b71799e79 --- /dev/null +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsProducerGPU.cc @@ -0,0 +1,119 @@ +#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/HostAllocator.h" + +#include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h" + +#include "EcalUncalibRecHitPhase2WeightsAlgoGPU.h" +#include "DeclsForKernelsPhase2.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 size_; +}; + +// constructor with initialisation of elements +EcalUncalibRecHitPhase2WeightsProducerGPU::EcalUncalibRecHitPhase2WeightsProducerGPU(const edm::ParameterSet &ps) + : // 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"))} {} + +void EcalUncalibRecHitPhase2WeightsProducerGPU::fillDescriptions(edm::ConfigurationDescriptions &descriptions) { + 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, + -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); + + size_ = digis.size; + + // if no digis stop here + if (size_ == 0) + return; + + auto 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(size_, ctx.stream()); + + ecal::weights::phase2Weights(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 = size_; + + // put into the event + ctx.emplace(event, recHitsToken_, std::move(eventOutputDataGPU_.recHits)); +} + +DEFINE_FWK_MODULE(EcalUncalibRecHitPhase2WeightsProducerGPU); 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 diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py index 63a3f7711dde3..5b5f37d8905ea 100644 --- a/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py +++ b/RecoLocalCalo/EcalRecProducers/python/ecalUncalibRecHitPhase2_cff.py @@ -1,4 +1,55 @@ 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.ecalPhase2DigiToGPUProducer_cfi import ecalPhase2DigiToGPUProducer as _ecalPhase2DigiToGPUProducer +ecalPhase2DigiToGPUProducer = _ecalPhase2DigiToGPUProducer.clone() + +# gpu weights +from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitPhase2GPU_cfi import ecalUncalibRecHitPhase2GPU as _ecalUncalibRecHitPhase2GPU +ecalUncalibRecHitPhase2GPU = _ecalUncalibRecHitPhase2GPU.clone( + digisLabelEB = ('ecalPhase2DigiToGPUProducer', 'ebDigis') +) + +# copy the uncalibrated rechits from GPU to CPU +from RecoLocalCalo.EcalRecProducers.ecalCPUUncalibRecHitProducer_cfi import ecalCPUUncalibRecHitProducer as _ecalCPUUncalibRecHitProducer +ecalUncalibRecHitSoA = _ecalCPUUncalibRecHitProducer.clone( + recHitsInLabelEB = ('ecalUncalibRecHitPhase2GPU', 'EcalUncalibRecHitsEB'), + isPhase2 = True, + recHitsInLabelEE = None, # remove unneeded Phase1 parameters + recHitsOutLabelEE = None +) + + +from RecoLocalCalo.EcalRecProducers.ecalUncalibRecHitConvertGPU2CPUFormat_cfi import ecalUncalibRecHitConvertGPU2CPUFormat as _ecalUncalibRecHitConvertGPU2CPUFormat +gpu.toModify(ecalUncalibRecHitPhase2, + 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 + ecalPhase2DigiToGPUProducer, + # ECAL weights running on GPU + ecalUncalibRecHitPhase2GPU, + # copy the uncalibrated rechits from GPU to CPU + ecalUncalibRecHitSoA, + # ECAL multifit running on CPU, or convert the uncalibrated rechits from SoA to legacy format + ecalUncalibRecHitPhase2, +)) 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 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