From 1148c5637eb0258745b13952810a81c6f7daae76 Mon Sep 17 00:00:00 2001 From: Andrea Massironi Date: Tue, 21 Jul 2020 15:48:57 +0200 Subject: [PATCH] Configure the number of ECAL barrel and endcap channels separately (#517) Fix memory allocation issues. Apply come code clean up: - remove outdated comments; - replace MYMALLOC macro with a lambda; - reuse named values from EcalDataFrame. --- .../EcalRawToDigi/plugins/DeclsForKernels.h | 18 +- .../plugins/EcalCPUDigisProducer.cc | 6 +- .../EcalRawToDigi/plugins/EcalRawToDigiGPU.cc | 6 +- .../plugins/DeclsForKernels.h | 163 ++++++++---------- .../plugins/EcalRecHitProducerGPU.cc | 12 +- .../plugins/EcalUncalibRecHitProducerGPU.cc | 12 +- .../python/ecalRecHitGPU_cfi.py | 4 +- 7 files changed, 104 insertions(+), 117 deletions(-) diff --git a/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h b/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h index ddb347dc21211..343b329165f55 100644 --- a/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h +++ b/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h @@ -3,6 +3,8 @@ #include +#include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" +#include "DataFormats/EcalDigi/interface/EcalDataFrame.h" #include "EventFilter/EcalRawToDigi/interface/DCCRawDataDefinitions.h" #include "EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h" #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -10,8 +12,6 @@ #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" -#include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" - namespace ecal { namespace raw { @@ -26,7 +26,8 @@ namespace ecal { }; struct ConfigurationParameters { - uint32_t maxChannels; + uint32_t maxChannelsEE; + uint32_t maxChannelsEB; }; struct OutputDataCPU { @@ -37,14 +38,11 @@ namespace ecal { struct OutputDataGPU { DigisCollection<::calo::common::DevStoragePolicy> digisEB, digisEE; - // FIXME: we should separate max channels parameter for eb and ee - // FIXME: replace hardcoded values void allocate(ConfigurationParameters const &config, cudaStream_t cudaStream) { - digisEB.data = cms::cuda::make_device_unique(config.maxChannels, cudaStream); - digisEE.data = cms::cuda::make_device_unique(config.maxChannels, cudaStream); - - digisEB.ids = cms::cuda::make_device_unique(config.maxChannels, cudaStream); - digisEE.ids = cms::cuda::make_device_unique(config.maxChannels, cudaStream); + digisEB.data = cms::cuda::make_device_unique(config.maxChannelsEB * EcalDataFrame::MAXSAMPLES, cudaStream); + digisEE.data = cms::cuda::make_device_unique(config.maxChannelsEE * EcalDataFrame::MAXSAMPLES, cudaStream); + digisEB.ids = cms::cuda::make_device_unique(config.maxChannelsEB, cudaStream); + digisEE.ids = cms::cuda::make_device_unique(config.maxChannelsEE, cudaStream); } }; diff --git a/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc b/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc index e39963c98110b..e3b283d3e91ec 100644 --- a/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc +++ b/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc @@ -3,6 +3,7 @@ #include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" #include "CondFormats/DataRecord/interface/EcalMappingElectronicsRcd.h" #include "DataFormats/EcalDetId/interface/EcalDetIdCollections.h" +#include "DataFormats/EcalDigi/interface/EcalDataFrame.h" #include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" #include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" #include "DataFormats/FEDRawData/interface/FEDRawDataCollection.h" @@ -134,11 +135,10 @@ void EcalCPUDigisProducer::acquire(edm::Event const& event, auto const& eedigis = ctx.get(eedigisProduct); // resize tmp buffers - // FIXME remove hardcoded values + dataebtmp.resize(ebdigis.size * EcalDataFrame::MAXSAMPLES); + dataeetmp.resize(eedigis.size * EcalDataFrame::MAXSAMPLES); idsebtmp.resize(ebdigis.size); - dataebtmp.resize(ebdigis.size * 10); idseetmp.resize(eedigis.size); - dataeetmp.resize(eedigis.size * 10); // enqeue transfers cudaCheck(cudaMemcpyAsync( diff --git a/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc b/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc index dd8e9992b220d..260567faf0a4b 100644 --- a/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc +++ b/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc @@ -48,7 +48,8 @@ void EcalRawToDigiGPU::fillDescriptions(edm::ConfigurationDescriptions& confDesc for (uint32_t i = 0; i < 54; ++i) feds[i] = i + 601; desc.add>("FEDs", feds); - desc.add("maxChannels", 20000); + desc.add("maxChannelsEB", 61200); + desc.add("maxChannelsEE", 14648); desc.add("digisLabelEB", "ebDigisGPU"); desc.add("digisLabelEE", "eeDigisGPU"); @@ -61,7 +62,8 @@ EcalRawToDigiGPU::EcalRawToDigiGPU(const edm::ParameterSet& ps) digisEBToken_{produces(ps.getParameter("digisLabelEB"))}, digisEEToken_{produces(ps.getParameter("digisLabelEE"))}, fedsToUnpack_{ps.getParameter>("FEDs")} { - config_.maxChannels = ps.getParameter("maxChannels"); + config_.maxChannelsEB = ps.getParameter("maxChannelsEB"); + config_.maxChannelsEE = ps.getParameter("maxChannelsEE"); } EcalRawToDigiGPU::~EcalRawToDigiGPU() {} diff --git a/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h index b6ceb52e3c41a..aa4f88c58c4e8 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h @@ -91,41 +91,43 @@ namespace ecal { bool shouldRunTimingComputation; - uint32_t maxNumberHits; + uint32_t maxNumberHitsEB; + uint32_t maxNumberHitsEE; }; struct EventOutputDataGPU { UncalibratedRecHit<::calo::common::DevStoragePolicy> recHitsEB, recHitsEE; void allocate(ConfigurationParameters const& configParameters, cudaStream_t cudaStream) { - auto const size = configParameters.maxNumberHits; + auto const sizeEB = configParameters.maxNumberHitsEB; recHitsEB.amplitudesAll = - cms::cuda::make_device_unique(size * EcalDataFrame::MAXSAMPLES, cudaStream); - recHitsEB.amplitude = cms::cuda::make_device_unique(size, cudaStream); - recHitsEB.chi2 = cms::cuda::make_device_unique(size, cudaStream); - recHitsEB.pedestal = cms::cuda::make_device_unique(size, cudaStream); + cms::cuda::make_device_unique(sizeEB * EcalDataFrame::MAXSAMPLES, cudaStream); + recHitsEB.amplitude = cms::cuda::make_device_unique(sizeEB, cudaStream); + recHitsEB.chi2 = cms::cuda::make_device_unique(sizeEB, cudaStream); + recHitsEB.pedestal = cms::cuda::make_device_unique(sizeEB, cudaStream); if (configParameters.shouldRunTimingComputation) { - recHitsEB.jitter = cms::cuda::make_device_unique(size, cudaStream); - recHitsEB.jitterError = cms::cuda::make_device_unique(size, cudaStream); + recHitsEB.jitter = cms::cuda::make_device_unique(sizeEB, cudaStream); + recHitsEB.jitterError = cms::cuda::make_device_unique(sizeEB, cudaStream); } - recHitsEB.did = cms::cuda::make_device_unique(size, cudaStream); - recHitsEB.flags = cms::cuda::make_device_unique(size, cudaStream); + recHitsEB.did = cms::cuda::make_device_unique(sizeEB, cudaStream); + recHitsEB.flags = cms::cuda::make_device_unique(sizeEB, cudaStream); + auto const sizeEE = configParameters.maxNumberHitsEE; recHitsEE.amplitudesAll = - cms::cuda::make_device_unique(size * EcalDataFrame::MAXSAMPLES, cudaStream); - recHitsEE.amplitude = cms::cuda::make_device_unique(size, cudaStream); - recHitsEE.chi2 = cms::cuda::make_device_unique(size, cudaStream); - recHitsEE.pedestal = cms::cuda::make_device_unique(size, cudaStream); + cms::cuda::make_device_unique(sizeEE * EcalDataFrame::MAXSAMPLES, cudaStream); + recHitsEE.amplitude = cms::cuda::make_device_unique(sizeEE, cudaStream); + recHitsEE.chi2 = cms::cuda::make_device_unique(sizeEE, cudaStream); + recHitsEE.pedestal = cms::cuda::make_device_unique(sizeEE, cudaStream); if (configParameters.shouldRunTimingComputation) { - recHitsEE.jitter = cms::cuda::make_device_unique(size, cudaStream); - recHitsEE.jitterError = cms::cuda::make_device_unique(size, cudaStream); + recHitsEE.jitter = cms::cuda::make_device_unique(sizeEE, cudaStream); + recHitsEE.jitterError = cms::cuda::make_device_unique(sizeEE, cudaStream); } - recHitsEE.did = cms::cuda::make_device_unique(size, cudaStream); - recHitsEE.flags = cms::cuda::make_device_unique(size, cudaStream); + recHitsEE.did = cms::cuda::make_device_unique(sizeEE, cudaStream); + recHitsEE.flags = cms::cuda::make_device_unique(sizeEE, cudaStream); } }; @@ -169,64 +171,43 @@ namespace ecal { constexpr auto smlength = getLength(); constexpr auto pmlength = getLength(); constexpr auto bxvlength = getLength(); - auto const size = configParameters.maxNumberHits; - -#define MYMALLOC(var, size) var = cms::cuda::make_device_unique(size, cudaStream) - MYMALLOC(samples, size * svlength); - //cudaCheck(cudaMalloc((void**)&samples, size * sizeof(SampleVector))); - MYMALLOC(gainsNoise, size * sgvlength); - //cudaCheck(cudaMalloc((void**)&gainsNoise, size * sizeof(SampleGainVector))); - - MYMALLOC(noisecov, size * smlength); - //cudaCheck(cudaMalloc((void**)&noisecov, size * sizeof(SampleMatrix))); - MYMALLOC(pulse_matrix, size * pmlength); - //cudaCheck(cudaMalloc((void**)&pulse_matrix, size * sizeof(PulseMatrixType))); - MYMALLOC(activeBXs, size * bxvlength); - //cudaCheck(cudaMalloc((void**)&activeBXs, size * sizeof(BXVectorType))); - MYMALLOC(acState, size); - //cudaCheck(cudaMalloc((void**)&acState, size * sizeof(char))); - - MYMALLOC(hasSwitchToGain6, size); - //cudaCheck(cudaMalloc((void**)&hasSwitchToGain6, size * sizeof(bool))); - MYMALLOC(hasSwitchToGain1, size); - //cudaCheck(cudaMalloc((void**)&hasSwitchToGain1, size * sizeof(bool))); - MYMALLOC(isSaturated, size); - //cudaCheck(cudaMalloc((void**)&isSaturated, size * sizeof(bool))); + auto const size = configParameters.maxNumberHitsEB + configParameters.maxNumberHitsEE; + + auto alloc = [cudaStream](auto& var, uint32_t size) { + using element_type = typename std::remove_reference_t::element_type; + var = cms::cuda::make_device_unique(size, cudaStream); + }; + + alloc(samples, size * svlength); + alloc(gainsNoise, size * sgvlength); + + alloc(noisecov, size * smlength); + alloc(pulse_matrix, size * pmlength); + alloc(activeBXs, size * bxvlength); + alloc(acState, size); + + alloc(hasSwitchToGain6, size); + alloc(hasSwitchToGain1, size); + alloc(isSaturated, size); if (configParameters.shouldRunTimingComputation) { - MYMALLOC(sample_values, size * svlength); - //cudaCheck(cudaMalloc((void**)&sample_values, size * sizeof(SampleVector))); - MYMALLOC(sample_value_errors, size * svlength); - //cudaCheck(cudaMalloc((void**)&sample_value_errors, size * sizeof(SampleVector))); - MYMALLOC(useless_sample_values, size * EcalDataFrame::MAXSAMPLES); - //cudaCheck(cudaMalloc((void**)&useless_sample_values, size * sizeof(bool) * EcalDataFrame::MAXSAMPLES)); - MYMALLOC(chi2sNullHypot, size); - //cudaCheck(cudaMalloc((void**)&chi2sNullHypot, size * sizeof(SampleVector::Scalar))); - MYMALLOC(sum0sNullHypot, size); - //cudaCheck(cudaMalloc((void**)&sum0sNullHypot, size * sizeof(SampleVector::Scalar))); - MYMALLOC(sumAAsNullHypot, size); - //cudaCheck(cudaMalloc((void**)&sumAAsNullHypot, size * sizeof(SampleVector::Scalar))); - MYMALLOC(pedestal_nums, size); - //cudaCheck(cudaMalloc((void**)&pedestal_nums, size * sizeof(char))); - - MYMALLOC(tMaxAlphaBetas, size); - //cudaCheck(cudaMalloc((void**)&tMaxAlphaBetas, size * sizeof(SampleVector::Scalar))); - MYMALLOC(tMaxErrorAlphaBetas, size); - //cudaCheck(cudaMalloc((void**)&tMaxErrorAlphaBetas, size * sizeof(SampleVector::Scalar))); - MYMALLOC(accTimeMax, size); - //cudaCheck(cudaMalloc((void**)&accTimeMax, size * sizeof(SampleVector::Scalar))); - MYMALLOC(accTimeWgt, size); - //cudaCheck(cudaMalloc((void**)&accTimeWgt, size * sizeof(SampleVector::Scalar))); - MYMALLOC(ampMaxAlphaBeta, size); - //cudaCheck(cudaMalloc((void**)&MaxAlphaBeta, size * sizeof(SampleVector::Scalar))); - MYMALLOC(ampMaxError, size); - //cudaCheck(cudaMalloc((void**)&MaxError, size * sizeof(SampleVector::Scalar))); - MYMALLOC(timeMax, size); - //cudaCheck(cudaMalloc((void**)&timeMax, size * sizeof(SampleVector::Scalar))); - MYMALLOC(timeError, size); - //cudaCheck(cudaMalloc((void**)&timeError, size * sizeof(SampleVector::Scalar))); - MYMALLOC(tcState, size); - //cudaCheck(cudaMalloc((void**)&tcState, size * sizeof(TimeComputationState))); + alloc(sample_values, size * svlength); + alloc(sample_value_errors, size * svlength); + alloc(useless_sample_values, size * EcalDataFrame::MAXSAMPLES); + alloc(chi2sNullHypot, size); + alloc(sum0sNullHypot, size); + alloc(sumAAsNullHypot, size); + alloc(pedestal_nums, size); + + alloc(tMaxAlphaBetas, size); + alloc(tMaxErrorAlphaBetas, size); + alloc(accTimeMax, size); + alloc(accTimeWgt, size); + alloc(ampMaxAlphaBeta, size); + alloc(ampMaxError, size); + alloc(timeMax, size); + alloc(timeError, size); + alloc(tcState, size); } } }; @@ -295,8 +276,9 @@ namespace ecal { uint32_t expanded_v_DB_reco_flagsSize; uint32_t flagmask; - uint32_t maxNumberHits; - + uint32_t maxNumberHitsEB; + uint32_t maxNumberHitsEE; + // // bool shouldRunTimingComputation; }; @@ -307,20 +289,21 @@ namespace ecal { void allocate(ConfigurationParameters const& configParameters, cudaStream_t cudaStream) { // void allocate(uint32_t size) { //---- configParameters -> needed only to decide if to save the timing information or not - auto const size = configParameters.maxNumberHits; - recHitsEB.energy = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(size, cudaStream); - recHitsEB.time = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(size, cudaStream); - recHitsEB.chi2 = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(size, cudaStream); - recHitsEB.flagBits = cms::cuda::make_device_unique(size, cudaStream); - recHitsEB.extra = cms::cuda::make_device_unique(size, cudaStream); - recHitsEB.did = cms::cuda::make_device_unique(size, cudaStream); - - recHitsEE.energy = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(size, cudaStream); - recHitsEE.time = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(size, cudaStream); - recHitsEE.chi2 = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(size, cudaStream); - recHitsEE.flagBits = cms::cuda::make_device_unique(size, cudaStream); - recHitsEE.extra = cms::cuda::make_device_unique(size, cudaStream); - recHitsEE.did = cms::cuda::make_device_unique(size, cudaStream); + auto const sizeEB = configParameters.maxNumberHitsEB; + recHitsEB.energy = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(sizeEB, cudaStream); + recHitsEB.time = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(sizeEB, cudaStream); + recHitsEB.chi2 = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(sizeEB, cudaStream); + recHitsEB.flagBits = cms::cuda::make_device_unique(sizeEB, cudaStream); + recHitsEB.extra = cms::cuda::make_device_unique(sizeEB, cudaStream); + recHitsEB.did = cms::cuda::make_device_unique(sizeEB, cudaStream); + + auto const sizeEE = configParameters.maxNumberHitsEE; + recHitsEE.energy = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(sizeEE, cudaStream); + recHitsEE.time = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(sizeEE, cudaStream); + recHitsEE.chi2 = cms::cuda::make_device_unique<::ecal::reco::StorageScalarType[]>(sizeEE, cudaStream); + recHitsEE.flagBits = cms::cuda::make_device_unique(sizeEE, cudaStream); + recHitsEE.extra = cms::cuda::make_device_unique(sizeEE, cudaStream); + recHitsEE.did = cms::cuda::make_device_unique(sizeEE, cudaStream); } }; diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc index 8000f8c39b773..f34925fe87ad8 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc @@ -100,7 +100,8 @@ void EcalRecHitProducerGPU::fillDescriptions(edm::ConfigurationDescriptions& con desc.add("EBLaserMAX", 30.0); desc.add("EELaserMAX", 30.0); - desc.add("maxNumberHits", 20000); + desc.add("maxNumberHitsEB", 61200); + desc.add("maxNumberHitsEE", 14648); } EcalRecHitProducerGPU::EcalRecHitProducerGPU(const edm::ParameterSet& ps) { @@ -121,8 +122,9 @@ EcalRecHitProducerGPU::EcalRecHitProducerGPU(const edm::ParameterSet& ps) { configParameters_.EELaserMAX = ps.getParameter("EELaserMAX"); // max number of digis to allocate for - configParameters_.maxNumberHits = ps.getParameter("maxNumberHits"); - + configParameters_.maxNumberHitsEB = ps.getParameter("maxNumberHitsEB"); + configParameters_.maxNumberHitsEE = ps.getParameter("maxNumberHitsEE"); + flagmask_ = 0; flagmask_ |= 0x1 << EcalRecHit::kNeighboursRecovered; flagmask_ |= 0x1 << EcalRecHit::kTowerRecovered; @@ -163,8 +165,8 @@ void EcalRecHitProducerGPU::acquire(edm::Event const& event, nee_ = eeUncalibRecHits.size; // std::cout << " [EcalRecHitProducerGPU::acquire] neb_:nee_ = " << neb_ << " : " << nee_ << std::endl; - if ((neb_ + nee_) > configParameters_.maxNumberHits) { - edm::LogError("EcalRecHitProducerGPU") << "max number of channels exceeded. See options 'maxNumberHits' "; + if ((neb_ > configParameters_.maxNumberHitsEB) || (nee_ > configParameters_.maxNumberHitsEE)) { + edm::LogError("EcalRecHitProducerGPU") << "max number of channels exceeded. See options 'maxNumberHitsEB and maxNumberHitsEE' "; } int nchannelsEB = ebUncalibRecHits.size; // --> offsetForInput, first EB and then EE diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc index 0e29fc0c3edd6..fed87e073eef8 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc @@ -99,7 +99,8 @@ void EcalUncalibRecHitProducerGPU::fillDescriptions(edm::ConfigurationDescriptio desc.add("outOfTimeThresholdGain61mEE", 1000); desc.add("amplitudeThresholdEB", 10); desc.add("amplitudeThresholdEE", 10); - desc.add("maxNumberHits", 20000); //---- AM TEST + desc.add("maxNumberHitsEB", 61200); + desc.add("maxNumberHitsEE", 14648); desc.add>("kernelMinimizeThreads", {32, 1, 1}); // ---- default false or true? It was set to true, but at HLT it is false desc.add("shouldRunTimingComputation", false); @@ -145,8 +146,9 @@ EcalUncalibRecHitProducerGPU::EcalUncalibRecHitProducerGPU(const edm::ParameterS auto amplitudeThreshEE = ps.getParameter("amplitudeThresholdEE"); // max number of digis to allocate for - configParameters_.maxNumberHits = ps.getParameter("maxNumberHits"); - + configParameters_.maxNumberHitsEB = ps.getParameter("maxNumberHitsEB"); + configParameters_.maxNumberHitsEE = ps.getParameter("maxNumberHitsEE"); + // switch to run timing computation kernels configParameters_.shouldRunTimingComputation = ps.getParameter("shouldRunTimingComputation"); @@ -210,8 +212,8 @@ void EcalUncalibRecHitProducerGPU::acquire(edm::Event const& event, neb_ = ebDigis.size; nee_ = eeDigis.size; - if ((neb_ + nee_) > configParameters_.maxNumberHits) { - edm::LogError("EcalUncalibRecHitProducerGPU") << "max number of channels exceeded. See options 'maxNumberHits' "; + if ((neb_ > configParameters_.maxNumberHitsEB) || (nee_ > configParameters_.maxNumberHitsEE)) { + edm::LogError("EcalUncalibRecHitProducerGPU") << "max number of channels exceeded. See options 'maxNumberHitsEB and maxNumberHitsEE' "; } // conditions diff --git a/RecoLocalCalo/EcalRecProducers/python/ecalRecHitGPU_cfi.py b/RecoLocalCalo/EcalRecProducers/python/ecalRecHitGPU_cfi.py index 76299519b51dc..f5796430803eb 100644 --- a/RecoLocalCalo/EcalRecProducers/python/ecalRecHitGPU_cfi.py +++ b/RecoLocalCalo/EcalRecProducers/python/ecalRecHitGPU_cfi.py @@ -13,8 +13,8 @@ recHitsLabelEB = cms.string("EcalRecHitsEB"), recHitsLabelEE = cms.string("EcalRecHitsEE"), - maxNumberHits = cms.uint32(20000), # FIXME AM - + maxNumberHitsEB = cms.uint32(61200), + maxNumberHitsEE = cms.uint32(14648), #EErechitCollection = cms.string('EcalRecHitsEE'), #EEuncalibRecHitCollection = cms.InputTag("ecalMultiFitUncalibRecHit","EcalUncalibRecHitsEE"),