diff --git a/CondFormats/EcalObjects/BuildFile.xml b/CondFormats/EcalObjects/BuildFile.xml index aa87dacc6a043..11a0a25bec707 100644 --- a/CondFormats/EcalObjects/BuildFile.xml +++ b/CondFormats/EcalObjects/BuildFile.xml @@ -3,10 +3,13 @@ + + + diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h b/CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h similarity index 65% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h rename to CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h index a3f65d0b509fc..27a4236acd6e1 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h +++ b/CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalGainRatiosGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalGainRatiosGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalGainRatiosGPU_h +#define CondFormats_EcalObjects_interface_EcalGainRatiosGPU_h #include "CondFormats/EcalObjects/interface/EcalGainRatios.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,8 @@ class EcalGainRatiosGPU { public: struct Product { - ~Product(); - float *gain12Over6 = nullptr, *gain6Over1 = nullptr; + edm::propagate_const_array> gain12Over6; + edm::propagate_const_array> gain6Over1; }; #ifndef __CUDACC__ @@ -40,4 +42,4 @@ class EcalGainRatiosGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalGainRatiosGPU_h +#endif // CondFormats_EcalObjects_interface_EcalGainRatiosGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h b/CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h similarity index 60% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h rename to CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h index 4b5401ff0316f..35c172a2920af 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h +++ b/CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalIntercalibConstantsGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalIntercalibConstantsGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalIntercalibConstantsGPU_h +#define CondFormats_EcalObjects_interface_EcalIntercalibConstantsGPU_h #include "CondFormats/EcalObjects/interface/EcalIntercalibConstants.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,7 @@ class EcalIntercalibConstantsGPU { public: struct Product { - ~Product(); - float* values = nullptr; + edm::propagate_const_array> values; }; #ifndef __CUDACC__ @@ -27,17 +28,17 @@ class EcalIntercalibConstantsGPU { // TODO: do this centrally // get offset for hashes. equals number of barrel items - uint32_t getOffset() const { return valuesEB_.size(); } + uint32_t getOffset() const { return offset_; } // static std::string name() { return std::string{"ecalIntercalibConstantsGPU"}; } private: - std::vector const& valuesEB_; - std::vector const& valuesEE_; + std::vector> values_; + uint32_t offset_; cms::cuda::ESProduct product_; #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalIntercalibConstantsGPU_h +#endif // CondFormats_EcalObjects_interface_EcalIntercalibConstantsGPU_h diff --git a/CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h b/CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h new file mode 100644 index 0000000000000..f68c405470556 --- /dev/null +++ b/CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h @@ -0,0 +1,54 @@ +#ifndef CondFormats_EcalObjects_interface_EcalLaserAPDPNRatiosGPU_h +#define CondFormats_EcalObjects_interface_EcalLaserAPDPNRatiosGPU_h + +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatios.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" + +#ifndef __CUDACC__ +#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" +#include "HeterogeneousCore/CUDACore/interface/ESProduct.h" +#endif // __CUDACC__ + +class EcalLaserAPDPNRatiosGPU { +public: + struct Product { + edm::propagate_const_array> p1; + edm::propagate_const_array> p2; + edm::propagate_const_array> p3; + edm::propagate_const_array> t1; + edm::propagate_const_array> t2; + edm::propagate_const_array> t3; + }; + +#ifndef __CUDACC__ + + // + EcalLaserAPDPNRatiosGPU(EcalLaserAPDPNRatios const &); + + // will call dealloation for Product thru ~Product + ~EcalLaserAPDPNRatiosGPU() = default; + + // get device pointers + Product const &getProduct(cudaStream_t) const; + + // + static std::string name() { return std::string{"ecalLaserAPDPNRatiosGPU"}; } + +private: + // in the future, we need to arrange so to avoid this copy on the host + // store eb first then ee + std::vector> p1_; + std::vector> p2_; + std::vector> p3_; + + std::vector> t1_; + std::vector> t2_; + std::vector> t3_; + + cms::cuda::ESProduct product_; + +#endif // __CUDACC__ +}; + +#endif // CondFormats_EcalObjects_interface_EcalLaserAPDPNRatiosGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h b/CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h similarity index 60% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h rename to CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h index 985bfd9579f7c..2b9514a081c14 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h +++ b/CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAPDPNRatiosRefGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAPDPNRatiosRefGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalLaserAPDPNRatiosRefGPU_h +#define CondFormats_EcalObjects_interface_EcalLaserAPDPNRatiosRefGPU_h #include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRef.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,7 @@ class EcalLaserAPDPNRatiosRefGPU { public: struct Product { - ~Product(); - float* values = nullptr; + edm::propagate_const_array> values; }; #ifndef __CUDACC__ @@ -27,17 +28,17 @@ class EcalLaserAPDPNRatiosRefGPU { // TODO: do this centrally // get offset for hashes. equals number of barrel items - uint32_t getOffset() const { return valuesEB_.size(); } + uint32_t getOffset() const { return offset_; } + uint32_t offset_; // static std::string name() { return std::string{"ecalLaserAPDPNRatiosRefGPU"}; } private: - std::vector const& valuesEB_; - std::vector const& valuesEE_; + std::vector> values_; cms::cuda::ESProduct product_; #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAPDPNRatiosRefGPU_h +#endif // CondFormats_EcalObjects_interface_EcalLaserAPDPNRatiosRefGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h b/CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h similarity index 59% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h rename to CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h index 9dd05e9ee3c4d..50599a6afb43f 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h +++ b/CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAlphasGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAlphasGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalLaserAlphasGPU_h +#define CondFormats_EcalObjects_interface_EcalLaserAlphasGPU_h #include "CondFormats/EcalObjects/interface/EcalLaserAlphas.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,7 @@ class EcalLaserAlphasGPU { public: struct Product { - ~Product(); - float* values = nullptr; + edm::propagate_const_array> values; }; #ifndef __CUDACC__ @@ -27,17 +28,17 @@ class EcalLaserAlphasGPU { // TODO: do this centrally // get offset for hashes. equals number of barrel items - uint32_t getOffset() const { return valuesEB_.size(); } + uint32_t getOffset() const { return offset_; } // static std::string name() { return std::string{"ecalLaserAlphasGPU"}; } private: - std::vector const& valuesEB_; - std::vector const& valuesEE_; + std::vector> values_; + uint32_t offset_; cms::cuda::ESProduct product_; #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAlphasGPU_h +#endif // CondFormats_EcalObjects_interface_EcalLaserAlphasGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h b/CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h similarity index 60% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h rename to CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h index 343bdf1dd1afc..8513224071f9d 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h +++ b/CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalLinearCorrectionsGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalLinearCorrectionsGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalLinearCorrectionsGPU_h +#define CondFormats_EcalObjects_interface_EcalLinearCorrectionsGPU_h #include "CondFormats/EcalObjects/interface/EcalLinearCorrections.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,13 +13,12 @@ class EcalLinearCorrectionsGPU { public: struct Product { - ~Product(); - float *p1 = nullptr; - float *p2 = nullptr; - float *p3 = nullptr; - edm::TimeValue_t *t1 = nullptr; - edm::TimeValue_t *t2 = nullptr; - edm::TimeValue_t *t3 = nullptr; + edm::propagate_const_array> p1; + edm::propagate_const_array> p2; + edm::propagate_const_array> p3; + edm::propagate_const_array> t1; + edm::propagate_const_array> t2; + edm::propagate_const_array> t3; }; #ifndef __CUDACC__ @@ -50,4 +51,4 @@ class EcalLinearCorrectionsGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalLinearCorrectionsGPU_h +#endif // CondFormats_EcalObjects_interface_EcalLinearCorrectionsGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h b/CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h similarity index 57% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h rename to CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h index 56aa0579ff77f..399cfc14ada40 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h +++ b/CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h @@ -1,9 +1,11 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalMultifitParametersGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalMultifitParametersGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalMultifitParametersGPU_h +#define CondFormats_EcalObjects_interface_EcalMultifitParametersGPU_h #include #include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -13,8 +15,10 @@ class EcalMultifitParametersGPU { public: struct Product { - ~Product(); - double *amplitudeFitParametersEB, *amplitudeFitParametersEE, *timeFitParametersEB, *timeFitParametersEE; + edm::propagate_const_array> amplitudeFitParametersEB; + edm::propagate_const_array> amplitudeFitParametersEE; + edm::propagate_const_array> timeFitParametersEB; + edm::propagate_const_array> timeFitParametersEE; }; #ifndef __CUDACC__ @@ -36,4 +40,4 @@ class EcalMultifitParametersGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalMultifitParametersGPU_h +#endif // CondFormats_EcalObjects_interface_EcalMultifitParametersGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h b/CondFormats/EcalObjects/interface/EcalPedestalsGPU.h similarity index 59% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h rename to CondFormats/EcalObjects/interface/EcalPedestalsGPU.h index 5387c422ddd9e..80bc2d3a3fb80 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h +++ b/CondFormats/EcalObjects/interface/EcalPedestalsGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalPedestalsGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalPedestalsGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalPedestalsGPU_h +#define CondFormats_EcalObjects_interface_EcalPedestalsGPU_h #include "CondFormats/EcalObjects/interface/EcalPedestals.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,9 +13,12 @@ class EcalPedestalsGPU { public: struct Product { - ~Product(); - float *mean_x12 = nullptr, *mean_x6 = nullptr, *mean_x1 = nullptr; - float *rms_x12 = nullptr, *rms_x6 = nullptr, *rms_x1 = nullptr; + edm::propagate_const_array> mean_x12; + edm::propagate_const_array> mean_x6; + edm::propagate_const_array> mean_x1; + edm::propagate_const_array> rms_x12; + edm::propagate_const_array> rms_x6; + edm::propagate_const_array> rms_x1; }; #ifndef __CUDACC__ @@ -44,4 +49,4 @@ class EcalPedestalsGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalPedestalsGPU_h +#endif // CondFormats_EcalObjects_interface_EcalPedestalsGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h b/CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h similarity index 82% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h rename to CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h index 6c5a3d9b95e2e..56207d8cb4ab3 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h +++ b/CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h @@ -1,5 +1,5 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalPulseCovariancesGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalPulseCovariancesGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalPulseCovariancesGPU_h +#define CondFormats_EcalObjects_interface_EcalPulseCovariancesGPU_h #include "CondFormats/EcalObjects/interface/EcalPulseCovariances.h" @@ -37,4 +37,4 @@ class EcalPulseCovariancesGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalPulseCovariancesGPU_h +#endif // CondFormats_EcalObjects_interface_EcalPulseCovariancesGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h b/CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h similarity index 82% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h rename to CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h index 3edb2c9bcdfd3..0caf954409b18 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h +++ b/CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h @@ -1,5 +1,5 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalPulseShapesGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalPulseShapesGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalPulseShapesGPU_h +#define CondFormats_EcalObjects_interface_EcalPulseShapesGPU_h #include "CondFormats/EcalObjects/interface/EcalPulseShapes.h" @@ -37,4 +37,4 @@ class EcalPulseShapesGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalPulseShapesGPU_h +#endif // CondFormats_EcalObjects_interface_EcalPulseShapesGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalRecHitParametersGPU.h b/CondFormats/EcalObjects/interface/EcalRecHitParametersGPU.h similarity index 58% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalRecHitParametersGPU.h rename to CondFormats/EcalObjects/interface/EcalRecHitParametersGPU.h index c5d3dd0388d15..050fc51bc5036 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalRecHitParametersGPU.h +++ b/CondFormats/EcalObjects/interface/EcalRecHitParametersGPU.h @@ -1,9 +1,11 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalRecHitParametersGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalRecHitParametersGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalRecHitParametersGPU_h +#define CondFormats_EcalObjects_interface_EcalRecHitParametersGPU_h #include #include "FWCore/ParameterSet/interface/ParameterSet.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -13,9 +15,10 @@ class EcalRecHitParametersGPU { public: struct Product { - ~Product(); - int *ChannelStatusToBeExcluded, *expanded_v_DB_reco_flags; - uint32_t *expanded_Sizes_v_DB_reco_flags, *expanded_flagbit_v_DB_reco_flags; + edm::propagate_const_array> channelStatusToBeExcluded; + edm::propagate_const_array> expanded_v_DB_reco_flags; + edm::propagate_const_array> expanded_Sizes_v_DB_reco_flags; + edm::propagate_const_array> expanded_flagbit_v_DB_reco_flags; }; #ifndef __CUDACC__ @@ -28,14 +31,14 @@ class EcalRecHitParametersGPU { using intvec = std::reference_wrapper> const>; using uint32vec = std::reference_wrapper> const>; std::tuple getValues() const { - return {ChannelStatusToBeExcluded_, + return {channelStatusToBeExcluded_, expanded_v_DB_reco_flags_, expanded_Sizes_v_DB_reco_flags_, expanded_flagbit_v_DB_reco_flags_}; } private: - std::vector> ChannelStatusToBeExcluded_; + std::vector> channelStatusToBeExcluded_; std::vector> expanded_v_DB_reco_flags_; std::vector> expanded_Sizes_v_DB_reco_flags_, expanded_flagbit_v_DB_reco_flags_; @@ -44,4 +47,4 @@ class EcalRecHitParametersGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalRecHitParametersGPU_h +#endif // CondFormats_EcalObjects_interface_EcalRecHitParametersGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h b/CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h similarity index 68% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h rename to CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h index 7d4d3cc60fd5c..a4f2dc1e555b5 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h +++ b/CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalRechitADCToGeVConstantGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalRechitADCToGeVConstantGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalRechitADCToGeVConstantGPU_h +#define CondFormats_EcalObjects_interface_EcalRechitADCToGeVConstantGPU_h #include "CondFormats/EcalObjects/interface/EcalADCToGeVConstant.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,7 @@ class EcalRechitADCToGeVConstantGPU { public: struct Product { - ~Product(); - float* adc2gev = nullptr; + edm::propagate_const_array> adc2gev; }; #ifndef __CUDACC__ @@ -39,4 +40,4 @@ class EcalRechitADCToGeVConstantGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalRechitADCToGeVConstantGPU_h +#endif // CondFormats_EcalObjects_interface_EcalRechitADCToGeVConstantGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h b/CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h similarity index 68% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h rename to CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h index bab99ab656c2d..93e0f6a8dd7c0 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h +++ b/CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalRechitChannelStatusGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalRechitChannelStatusGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalRechitChannelStatusGPU_h +#define CondFormats_EcalObjects_interface_EcalRechitChannelStatusGPU_h #include "CondFormats/EcalObjects/interface/EcalChannelStatus.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,7 @@ class EcalRechitChannelStatusGPU { public: struct Product { - ~Product(); - uint16_t* status = nullptr; + edm::propagate_const_array> status; }; #ifndef __CUDACC__ @@ -39,4 +40,4 @@ class EcalRechitChannelStatusGPU { #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalRechitChannelStatusGPU_h +#endif // CondFormats_EcalObjects_interface_EcalRechitChannelStatusGPU_h diff --git a/CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h b/CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h new file mode 100644 index 0000000000000..dd6d48b16a6fa --- /dev/null +++ b/CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h @@ -0,0 +1,49 @@ +#ifndef CondFormats_EcalObjects_interface_EcalSamplesCorrelationGPU_h +#define CondFormats_EcalObjects_interface_EcalSamplesCorrelationGPU_h + +#include "CondFormats/EcalObjects/interface/EcalSamplesCorrelation.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" + +#ifndef __CUDACC__ +#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" +#include "HeterogeneousCore/CUDACore/interface/ESProduct.h" +#endif // __CUDACC__ + +class EcalSamplesCorrelationGPU { +public: + struct Product { + edm::propagate_const_array> EBG12SamplesCorrelation; + edm::propagate_const_array> EBG6SamplesCorrelation; + edm::propagate_const_array> EBG1SamplesCorrelation; + edm::propagate_const_array> EEG12SamplesCorrelation; + edm::propagate_const_array> EEG6SamplesCorrelation; + edm::propagate_const_array> EEG1SamplesCorrelation; + }; + +#ifndef __CUDACC__ + // rearrange pedestals + EcalSamplesCorrelationGPU(EcalSamplesCorrelation const&); + + // will call dealloation for Product thru ~Product + ~EcalSamplesCorrelationGPU() = default; + + // get device pointers + Product const& getProduct(cudaStream_t) const; + + // + static std::string name() { return std::string{"ecalSamplesCorrelationGPU"}; } + +private: + std::vector> EBG12SamplesCorrelation_; + std::vector> EBG6SamplesCorrelation_; + std::vector> EBG1SamplesCorrelation_; + std::vector> EEG12SamplesCorrelation_; + std::vector> EEG6SamplesCorrelation_; + std::vector> EEG1SamplesCorrelation_; + + cms::cuda::ESProduct product_; +#endif // __CUDACC__ +}; + +#endif // CondFormats_EcalObjects_interface_EcalSamplesCorrelationGPU_h diff --git a/CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h b/CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h new file mode 100644 index 0000000000000..086decb53f01b --- /dev/null +++ b/CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h @@ -0,0 +1,49 @@ +#ifndef CondFormats_EcalObjects_interface_EcalTimeBiasCorrectionsGPU_h +#define CondFormats_EcalObjects_interface_EcalTimeBiasCorrectionsGPU_h + +#include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrections.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" + +#ifndef __CUDACC__ +#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" +#include "HeterogeneousCore/CUDACore/interface/ESProduct.h" +#endif // __CUDACC__ + +class EcalTimeBiasCorrectionsGPU { +public: + struct Product { + edm::propagate_const_array> ebTimeCorrAmplitudeBins; + edm::propagate_const_array> ebTimeCorrShiftBins; + edm::propagate_const_array> eeTimeCorrAmplitudeBins; + edm::propagate_const_array> eeTimeCorrShiftBins; + int ebTimeCorrAmplitudeBinsSize, eeTimeCorrAmplitudeBinsSize; + }; + + // rearrange pedestals + EcalTimeBiasCorrectionsGPU(EcalTimeBiasCorrections const&); + +#ifndef __CUDACC__ + + // will call dealloation for Product thru ~Product + ~EcalTimeBiasCorrectionsGPU() = default; + + // get device pointers + Product const& getProduct(cudaStream_t) const; + + // + static std::string name() { return std::string{"ecalTimeBiasCorrectionsGPU"}; } +#endif // __CUDACC__ + +private: + std::vector> ebTimeCorrAmplitudeBins_; + std::vector> ebTimeCorrShiftBins_; + std::vector> eeTimeCorrAmplitudeBins_; + std::vector> eeTimeCorrShiftBins_; + +#ifndef __CUDACC__ + cms::cuda::ESProduct product_; +#endif // __CUDACC__ +}; + +#endif // CondFormats_EcalObjects_interface_EcalTimeBiasCorrectionsGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h b/CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h similarity index 60% rename from RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h rename to CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h index 823334d433cc2..410e7699bdf7b 100644 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h +++ b/CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h @@ -1,7 +1,9 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalTimeCalibConstantsGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalTimeCalibConstantsGPU_h +#ifndef CondFormats_EcalObjects_interface_EcalTimeCalibConstantsGPU_h +#define CondFormats_EcalObjects_interface_EcalTimeCalibConstantsGPU_h #include "CondFormats/EcalObjects/interface/EcalTimeCalibConstants.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -11,8 +13,7 @@ class EcalTimeCalibConstantsGPU { public: struct Product { - ~Product(); - float* values = nullptr; + edm::propagate_const_array> values; }; #ifndef __CUDACC__ @@ -27,17 +28,17 @@ class EcalTimeCalibConstantsGPU { // TODO: do this centrally // get offset for hashes. equals number of barrel items - uint32_t getOffset() const { return valuesEB_.size(); } + uint32_t getOffset() const { return offset_; } // static std::string name() { return std::string{"ecalTimeCalibConstantsGPU"}; } private: - std::vector const& valuesEB_; - std::vector const& valuesEE_; + std::vector> values_; + uint32_t offset_; cms::cuda::ESProduct product_; #endif // __CUDACC__ }; -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalTimeCalibConstantsGPU_h +#endif // CondFormats_EcalObjects_interface_EcalTimeCalibConstantsGPU_h diff --git a/EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h b/CondFormats/EcalObjects/interface/ElectronicsMappingGPU.h similarity index 71% rename from EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h rename to CondFormats/EcalObjects/interface/ElectronicsMappingGPU.h index 004821afe3ed8..4c0cb30222bb9 100644 --- a/EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h +++ b/CondFormats/EcalObjects/interface/ElectronicsMappingGPU.h @@ -1,7 +1,9 @@ -#ifndef EventFilter_EcalRawToDigi_interface_ElectronicsMappingGPU_h -#define EventFilter_EcalRawToDigi_interface_ElectronicsMappingGPU_h +#ifndef CondFormats_EcalObjects_interface_ElectronicsMappingGPU_h +#define CondFormats_EcalObjects_interface_ElectronicsMappingGPU_h #include "CondFormats/EcalObjects/interface/EcalMappingElectronics.h" +#include "FWCore/Utilities/interface/propagate_const_array.h" +#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #ifndef __CUDACC__ #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" @@ -14,8 +16,7 @@ namespace ecal { class ElectronicsMappingGPU { public: struct Product { - ~Product(); - uint32_t* eid2did; + edm::propagate_const_array> eid2did; }; #ifndef __CUDACC__ @@ -44,4 +45,4 @@ namespace ecal { } // namespace raw } // namespace ecal -#endif // EventFilter_EcalRawToDigi_interface_ElectronicsMappingGPU_h +#endif // CondFormats_EcalObjects_interface_ElectronicsMappingGPU_h diff --git a/CondFormats/EcalObjects/src/EcalGainRatiosGPU.cc b/CondFormats/EcalObjects/src/EcalGainRatiosGPU.cc new file mode 100644 index 0000000000000..9e509430c4276 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalGainRatiosGPU.cc @@ -0,0 +1,38 @@ +#include "CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalGainRatiosGPU::EcalGainRatiosGPU(EcalGainRatios const& values) + : gain12Over6_(values.size()), gain6Over1_(values.size()) { + // fill in eb + auto const& barrelValues = values.barrelItems(); + for (unsigned int i = 0; i < barrelValues.size(); i++) { + gain12Over6_[i] = barrelValues[i].gain12Over6(); + gain6Over1_[i] = barrelValues[i].gain6Over1(); + } + + // fill in ee + auto const& endcapValues = values.endcapItems(); + auto const offset = barrelValues.size(); + for (unsigned int i = 0; i < endcapValues.size(); i++) { + gain12Over6_[offset + i] = endcapValues[i].gain12Over6(); + gain6Over1_[offset + i] = endcapValues[i].gain6Over1(); + } +} + +EcalGainRatiosGPU::Product const& EcalGainRatiosGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalGainRatiosGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.gain12Over6 = cms::cuda::make_device_unique(gain12Over6_.size(), cudaStream); + product.gain6Over1 = cms::cuda::make_device_unique(gain6Over1_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.gain12Over6, gain12Over6_, cudaStream); + cms::cuda::copyAsync(product.gain6Over1, gain6Over1_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalGainRatiosGPU); diff --git a/CondFormats/EcalObjects/src/EcalIntercalibConstantsGPU.cc b/CondFormats/EcalObjects/src/EcalIntercalibConstantsGPU.cc new file mode 100644 index 0000000000000..c1e462a8c5a75 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalIntercalibConstantsGPU.cc @@ -0,0 +1,24 @@ +#include "CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalIntercalibConstantsGPU::EcalIntercalibConstantsGPU(EcalIntercalibConstants const& values) { + values_.reserve(values.size()); + std::copy(values.begin(), values.end(), values_.begin()); + offset_ = values.barrelItems().size(); +} + +EcalIntercalibConstantsGPU::Product const& EcalIntercalibConstantsGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalIntercalibConstantsGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.values = cms::cuda::make_device_unique(values_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.values, values_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalIntercalibConstantsGPU); diff --git a/CondFormats/EcalObjects/src/EcalLaserAPDPNRatiosGPU.cc b/CondFormats/EcalObjects/src/EcalLaserAPDPNRatiosGPU.cc new file mode 100644 index 0000000000000..3a74f029ec185 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalLaserAPDPNRatiosGPU.cc @@ -0,0 +1,61 @@ +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalLaserAPDPNRatiosGPU::EcalLaserAPDPNRatiosGPU(EcalLaserAPDPNRatios const& values) + : p1_(values.getLaserMap().size()), + p2_(values.getLaserMap().size()), + p3_(values.getLaserMap().size()), + t1_(values.getTimeMap().size()), + t2_(values.getTimeMap().size()), + t3_(values.getTimeMap().size()) { + // fill in eb + // auto const& barrelValues = values.barrelItems(); + for (unsigned int i = 0; i < values.getLaserMap().barrelItems().size(); i++) { + p1_[i] = values.getLaserMap().barrelItems()[i].p1; + p2_[i] = values.getLaserMap().barrelItems()[i].p2; + p3_[i] = values.getLaserMap().barrelItems()[i].p3; + } + + // fill in ee + // auto const& endcapValues = values.endcapItems(); + auto const offset_laser = values.getLaserMap().barrelItems().size(); + for (unsigned int i = 0; i < values.getLaserMap().endcapItems().size(); i++) { + p1_[offset_laser + i] = values.getLaserMap().endcapItems()[i].p1; + p2_[offset_laser + i] = values.getLaserMap().endcapItems()[i].p2; + p3_[offset_laser + i] = values.getLaserMap().endcapItems()[i].p3; + } + + // Time is a simple std::vector + // typedef std::vector EcalLaserTimeStampMap; + for (unsigned int i = 0; i < values.getTimeMap().size(); i++) { + t1_[i] = values.getTimeMap()[i].t1.value(); + t2_[i] = values.getTimeMap()[i].t2.value(); + t3_[i] = values.getTimeMap()[i].t3.value(); + } +} + +EcalLaserAPDPNRatiosGPU::Product const& EcalLaserAPDPNRatiosGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalLaserAPDPNRatiosGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.p1 = cms::cuda::make_device_unique(p1_.size(), cudaStream); + product.p2 = cms::cuda::make_device_unique(p2_.size(), cudaStream); + product.p3 = cms::cuda::make_device_unique(p3_.size(), cudaStream); + product.t1 = cms::cuda::make_device_unique(t1_.size(), cudaStream); + product.t2 = cms::cuda::make_device_unique(t2_.size(), cudaStream); + product.t3 = cms::cuda::make_device_unique(t3_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.p1, p1_, cudaStream); + cms::cuda::copyAsync(product.p2, p2_, cudaStream); + cms::cuda::copyAsync(product.p3, p3_, cudaStream); + cms::cuda::copyAsync(product.t1, t1_, cudaStream); + cms::cuda::copyAsync(product.t2, t2_, cudaStream); + cms::cuda::copyAsync(product.t3, t3_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalLaserAPDPNRatiosGPU); diff --git a/CondFormats/EcalObjects/src/EcalLaserAPDPNRatiosRefGPU.cc b/CondFormats/EcalObjects/src/EcalLaserAPDPNRatiosRefGPU.cc new file mode 100644 index 0000000000000..ed06c6591597d --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalLaserAPDPNRatiosRefGPU.cc @@ -0,0 +1,24 @@ +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalLaserAPDPNRatiosRefGPU::EcalLaserAPDPNRatiosRefGPU(EcalLaserAPDPNRatiosRef const& values) { + values_.reserve(values.size()); + std::copy(values.begin(), values.end(), values_.begin()); + offset_ = values.barrelItems().size(); +} + +EcalLaserAPDPNRatiosRefGPU::Product const& EcalLaserAPDPNRatiosRefGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalLaserAPDPNRatiosRefGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.values = cms::cuda::make_device_unique(values_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.values, values_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalLaserAPDPNRatiosRefGPU); diff --git a/CondFormats/EcalObjects/src/EcalLaserAlphasGPU.cc b/CondFormats/EcalObjects/src/EcalLaserAlphasGPU.cc new file mode 100644 index 0000000000000..b16742f4964c8 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalLaserAlphasGPU.cc @@ -0,0 +1,24 @@ +#include "CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalLaserAlphasGPU::EcalLaserAlphasGPU(EcalLaserAlphas const& values) { + values_.reserve(values.size()); + std::copy(values.begin(), values.end(), values_.begin()); + offset_ = values.barrelItems().size(); +} + +EcalLaserAlphasGPU::Product const& EcalLaserAlphasGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalLaserAlphasGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.values = cms::cuda::make_device_unique(values_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.values, values_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalLaserAlphasGPU); diff --git a/CondFormats/EcalObjects/src/EcalLinearCorrectionsGPU.cc b/CondFormats/EcalObjects/src/EcalLinearCorrectionsGPU.cc new file mode 100644 index 0000000000000..b6b706d6b18e8 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalLinearCorrectionsGPU.cc @@ -0,0 +1,59 @@ +#include "CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalLinearCorrectionsGPU::EcalLinearCorrectionsGPU(EcalLinearCorrections const& values) + : p1_(values.getValueMap().size()), + p2_(values.getValueMap().size()), + p3_(values.getValueMap().size()), + t1_(values.getTimeMap().size()), + t2_(values.getTimeMap().size()), + t3_(values.getTimeMap().size()) { + // fill in eb + for (unsigned int i = 0; i < values.getValueMap().barrelItems().size(); i++) { + p1_[i] = values.getValueMap().barrelItems()[i].p1; + p2_[i] = values.getValueMap().barrelItems()[i].p2; + p3_[i] = values.getValueMap().barrelItems()[i].p3; + } + + // fill in ee + auto const offset_laser = values.getValueMap().barrelItems().size(); + for (unsigned int i = 0; i < values.getValueMap().endcapItems().size(); i++) { + p1_[offset_laser + i] = values.getValueMap().endcapItems()[i].p1; + p2_[offset_laser + i] = values.getValueMap().endcapItems()[i].p2; + p3_[offset_laser + i] = values.getValueMap().endcapItems()[i].p3; + } + + // Time is a simple std::vector + // typedef std::vector EcalLaserTimeStampMap; + for (unsigned int i = 0; i < values.getTimeMap().size(); i++) { + t1_[i] = values.getTimeMap()[i].t1.value(); + t2_[i] = values.getTimeMap()[i].t2.value(); + t3_[i] = values.getTimeMap()[i].t3.value(); + } +} + +EcalLinearCorrectionsGPU::Product const& EcalLinearCorrectionsGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalLinearCorrectionsGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.p1 = cms::cuda::make_device_unique(p1_.size(), cudaStream); + product.p2 = cms::cuda::make_device_unique(p2_.size(), cudaStream); + product.p3 = cms::cuda::make_device_unique(p3_.size(), cudaStream); + product.t1 = cms::cuda::make_device_unique(t1_.size(), cudaStream); + product.t2 = cms::cuda::make_device_unique(t2_.size(), cudaStream); + product.t3 = cms::cuda::make_device_unique(t3_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.p1, p1_, cudaStream); + cms::cuda::copyAsync(product.p2, p2_, cudaStream); + cms::cuda::copyAsync(product.p3, p3_, cudaStream); + cms::cuda::copyAsync(product.t1, t1_, cudaStream); + cms::cuda::copyAsync(product.t2, t2_, cudaStream); + cms::cuda::copyAsync(product.t3, t3_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalLinearCorrectionsGPU); diff --git a/CondFormats/EcalObjects/src/EcalMultifitParametersGPU.cc b/CondFormats/EcalObjects/src/EcalMultifitParametersGPU.cc new file mode 100644 index 0000000000000..204db7e0a58ab --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalMultifitParametersGPU.cc @@ -0,0 +1,42 @@ +#include "CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalMultifitParametersGPU::EcalMultifitParametersGPU(edm::ParameterSet const& ps) { + auto const& amplitudeFitParametersEB = ps.getParameter>("EBamplitudeFitParameters"); + auto const& amplitudeFitParametersEE = ps.getParameter>("EEamplitudeFitParameters"); + auto const& timeFitParametersEB = ps.getParameter>("EBtimeFitParameters"); + auto const& timeFitParametersEE = ps.getParameter>("EEtimeFitParameters"); + + amplitudeFitParametersEB_.resize(amplitudeFitParametersEB.size()); + amplitudeFitParametersEE_.resize(amplitudeFitParametersEE.size()); + timeFitParametersEB_.resize(timeFitParametersEB.size()); + timeFitParametersEE_.resize(timeFitParametersEE.size()); + + std::copy(amplitudeFitParametersEB.begin(), amplitudeFitParametersEB.end(), amplitudeFitParametersEB_.begin()); + std::copy(amplitudeFitParametersEE.begin(), amplitudeFitParametersEE.end(), amplitudeFitParametersEE_.begin()); + std::copy(timeFitParametersEB.begin(), timeFitParametersEB.end(), timeFitParametersEB_.begin()); + std::copy(timeFitParametersEE.begin(), timeFitParametersEE.end(), timeFitParametersEE_.begin()); +} + +EcalMultifitParametersGPU::Product const& EcalMultifitParametersGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalMultifitParametersGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.amplitudeFitParametersEB = + cms::cuda::make_device_unique(amplitudeFitParametersEB_.size(), cudaStream); + product.amplitudeFitParametersEE = + cms::cuda::make_device_unique(amplitudeFitParametersEE_.size(), cudaStream); + product.timeFitParametersEB = cms::cuda::make_device_unique(timeFitParametersEB_.size(), cudaStream); + product.timeFitParametersEE = cms::cuda::make_device_unique(timeFitParametersEE_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.amplitudeFitParametersEB, amplitudeFitParametersEB_, cudaStream); + cms::cuda::copyAsync(product.amplitudeFitParametersEE, amplitudeFitParametersEE_, cudaStream); + cms::cuda::copyAsync(product.timeFitParametersEB, timeFitParametersEB_, cudaStream); + cms::cuda::copyAsync(product.timeFitParametersEE, timeFitParametersEE_, cudaStream); + }); + return product; +} + +TYPELOOKUP_DATA_REG(EcalMultifitParametersGPU); diff --git a/CondFormats/EcalObjects/src/EcalPedestalsGPU.cc b/CondFormats/EcalObjects/src/EcalPedestalsGPU.cc new file mode 100644 index 0000000000000..2528dab1665fb --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalPedestalsGPU.cc @@ -0,0 +1,59 @@ +#include "CondFormats/EcalObjects/interface/EcalPedestalsGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalPedestalsGPU::EcalPedestalsGPU(EcalPedestals const& pedestals) + : mean_x12_(pedestals.size()), + rms_x12_(pedestals.size()), + mean_x6_(pedestals.size()), + rms_x6_(pedestals.size()), + mean_x1_(pedestals.size()), + rms_x1_(pedestals.size()) { + // fill in eb + auto const& barrelValues = pedestals.barrelItems(); + for (unsigned int i = 0; i < barrelValues.size(); i++) { + mean_x12_[i] = barrelValues[i].mean_x12; + rms_x12_[i] = barrelValues[i].rms_x12; + mean_x6_[i] = barrelValues[i].mean_x6; + rms_x6_[i] = barrelValues[i].rms_x6; + mean_x1_[i] = barrelValues[i].mean_x1; + rms_x1_[i] = barrelValues[i].rms_x1; + } + + // fill in ee + auto const& endcapValues = pedestals.endcapItems(); + auto const offset = barrelValues.size(); + for (unsigned int i = 0; i < endcapValues.size(); i++) { + mean_x12_[offset + i] = endcapValues[i].mean_x12; + rms_x12_[offset + i] = endcapValues[i].rms_x12; + mean_x6_[offset + i] = endcapValues[i].mean_x6; + rms_x6_[offset + i] = endcapValues[i].rms_x6; + mean_x1_[offset + i] = endcapValues[i].mean_x1; + rms_x1_[offset + i] = endcapValues[i].rms_x1; + } +} + +EcalPedestalsGPU::Product const& EcalPedestalsGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalPedestalsGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.mean_x12 = cms::cuda::make_device_unique(mean_x12_.size(), cudaStream); + product.mean_x6 = cms::cuda::make_device_unique(mean_x6_.size(), cudaStream); + product.mean_x1 = cms::cuda::make_device_unique(mean_x1_.size(), cudaStream); + product.rms_x12 = cms::cuda::make_device_unique(rms_x12_.size(), cudaStream); + product.rms_x6 = cms::cuda::make_device_unique(rms_x6_.size(), cudaStream); + product.rms_x1 = cms::cuda::make_device_unique(rms_x1_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.mean_x12, mean_x12_, cudaStream); + cms::cuda::copyAsync(product.mean_x6, mean_x6_, cudaStream); + cms::cuda::copyAsync(product.mean_x1, mean_x1_, cudaStream); + cms::cuda::copyAsync(product.rms_x12, rms_x12_, cudaStream); + cms::cuda::copyAsync(product.rms_x6, rms_x6_, cudaStream); + cms::cuda::copyAsync(product.rms_x1, rms_x1_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalPedestalsGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalPulseCovariancesGPU.cc b/CondFormats/EcalObjects/src/EcalPulseCovariancesGPU.cc similarity index 95% rename from RecoLocalCalo/EcalRecAlgos/src/EcalPulseCovariancesGPU.cc rename to CondFormats/EcalObjects/src/EcalPulseCovariancesGPU.cc index bbeda99652e22..05586080bff6a 100644 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalPulseCovariancesGPU.cc +++ b/CondFormats/EcalObjects/src/EcalPulseCovariancesGPU.cc @@ -1,4 +1,4 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h" #include "FWCore/Utilities/interface/typelookup.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalPulseShapesGPU.cc b/CondFormats/EcalObjects/src/EcalPulseShapesGPU.cc similarity index 96% rename from RecoLocalCalo/EcalRecAlgos/src/EcalPulseShapesGPU.cc rename to CondFormats/EcalObjects/src/EcalPulseShapesGPU.cc index aee122a01627d..96b492f3ad335 100644 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalPulseShapesGPU.cc +++ b/CondFormats/EcalObjects/src/EcalPulseShapesGPU.cc @@ -1,4 +1,4 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h" #include "FWCore/Utilities/interface/typelookup.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" diff --git a/CondFormats/EcalObjects/src/EcalRecHitParametersGPU.cc b/CondFormats/EcalObjects/src/EcalRecHitParametersGPU.cc new file mode 100644 index 0000000000000..a74232bc93e8b --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalRecHitParametersGPU.cc @@ -0,0 +1,57 @@ +#include "CondFormats/EcalObjects/interface/EcalRecHitParametersGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" +#include "CommonTools/Utils/interface/StringToEnumValue.h" +#include "CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h" +#include "DataFormats/EcalRecHit/interface/EcalRecHit.h" + +EcalRecHitParametersGPU::EcalRecHitParametersGPU(edm::ParameterSet const& ps) { + auto const& channelStatusToBeExcluded = StringToEnumValue( + ps.getParameter>("ChannelStatusToBeExcluded")); + + channelStatusToBeExcluded_.resize(channelStatusToBeExcluded.size()); + std::copy(channelStatusToBeExcluded.begin(), channelStatusToBeExcluded.end(), channelStatusToBeExcluded_.begin()); + + // https://github.com/cms-sw/cmssw/blob/266e21cfc9eb409b093e4cf064f4c0a24c6ac293/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitWorkerSimple.cc + + // Translate string representation of flagsMapDBReco into enum values + const edm::ParameterSet& p = ps.getParameter("flagsMapDBReco"); + std::vector recoflagbitsStrings = p.getParameterNames(); + + for (unsigned int i = 0; i != recoflagbitsStrings.size(); ++i) { + EcalRecHit::Flags recoflagbit = (EcalRecHit::Flags)StringToEnumValue(recoflagbitsStrings[i]); + std::vector dbstatus_s = p.getParameter>(recoflagbitsStrings[i]); + for (unsigned int j = 0; j != dbstatus_s.size(); ++j) { + EcalChannelStatusCode::Code dbstatus = + (EcalChannelStatusCode::Code)StringToEnumValue(dbstatus_s[j]); + expanded_v_DB_reco_flags_.push_back(dbstatus); + } + + expanded_Sizes_v_DB_reco_flags_.push_back(dbstatus_s.size()); + expanded_flagbit_v_DB_reco_flags_.push_back(recoflagbit); + } +} + +EcalRecHitParametersGPU::Product const& EcalRecHitParametersGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalRecHitParametersGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.channelStatusToBeExcluded = + cms::cuda::make_device_unique(channelStatusToBeExcluded_.size(), cudaStream); + product.expanded_v_DB_reco_flags = + cms::cuda::make_device_unique(expanded_v_DB_reco_flags_.size(), cudaStream); + product.expanded_Sizes_v_DB_reco_flags = + cms::cuda::make_device_unique(expanded_Sizes_v_DB_reco_flags_.size(), cudaStream); + product.expanded_flagbit_v_DB_reco_flags = + cms::cuda::make_device_unique(expanded_flagbit_v_DB_reco_flags_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.channelStatusToBeExcluded, channelStatusToBeExcluded_, cudaStream); + cms::cuda::copyAsync(product.expanded_v_DB_reco_flags, expanded_v_DB_reco_flags_, cudaStream); + cms::cuda::copyAsync(product.expanded_Sizes_v_DB_reco_flags, expanded_Sizes_v_DB_reco_flags_, cudaStream); + cms::cuda::copyAsync(product.expanded_flagbit_v_DB_reco_flags, expanded_flagbit_v_DB_reco_flags_, cudaStream); + }); + return product; +} + +TYPELOOKUP_DATA_REG(EcalRecHitParametersGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalRechitADCToGeVConstantGPU.cc b/CondFormats/EcalObjects/src/EcalRechitADCToGeVConstantGPU.cc similarity index 50% rename from RecoLocalCalo/EcalRecAlgos/src/EcalRechitADCToGeVConstantGPU.cc rename to CondFormats/EcalObjects/src/EcalRechitADCToGeVConstantGPU.cc index 5f01068f95186..a71266d1615dc 100644 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalRechitADCToGeVConstantGPU.cc +++ b/CondFormats/EcalObjects/src/EcalRechitADCToGeVConstantGPU.cc @@ -1,7 +1,7 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h" #include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" EcalRechitADCToGeVConstantGPU::EcalRechitADCToGeVConstantGPU(EcalADCToGeVConstant const& values) : adc2gev_(2) // size is 2, one form EB and one for EE @@ -10,22 +10,13 @@ EcalRechitADCToGeVConstantGPU::EcalRechitADCToGeVConstantGPU(EcalADCToGeVConstan adc2gev_[1] = values.getEEValue(); } -EcalRechitADCToGeVConstantGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(adc2gev)); -} - EcalRechitADCToGeVConstantGPU::Product const& EcalRechitADCToGeVConstantGPU::getProduct(cudaStream_t cudaStream) const { auto const& product = product_.dataForCurrentDeviceAsync( cudaStream, [this](EcalRechitADCToGeVConstantGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.adc2gev, this->adc2gev_.size() * sizeof(float))); + // allocate + product.adc2gev = cms::cuda::make_device_unique(adc2gev_.size(), cudaStream); // transfer - cudaCheck(cudaMemcpyAsync(product.adc2gev, - this->adc2gev_.data(), - this->adc2gev_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); + cms::cuda::copyAsync(product.adc2gev, adc2gev_, cudaStream); }); return product; diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalRechitChannelStatusGPU.cc b/CondFormats/EcalObjects/src/EcalRechitChannelStatusGPU.cc similarity index 60% rename from RecoLocalCalo/EcalRecAlgos/src/EcalRechitChannelStatusGPU.cc rename to CondFormats/EcalObjects/src/EcalRechitChannelStatusGPU.cc index 1e6801fbd326a..2a0740828f6cf 100644 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalRechitChannelStatusGPU.cc +++ b/CondFormats/EcalObjects/src/EcalRechitChannelStatusGPU.cc @@ -1,7 +1,7 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h" #include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" EcalRechitChannelStatusGPU::EcalRechitChannelStatusGPU(EcalChannelStatus const& values) : status_(values.size()) { // fill in eb @@ -18,22 +18,13 @@ EcalRechitChannelStatusGPU::EcalRechitChannelStatusGPU(EcalChannelStatus const& } } -EcalRechitChannelStatusGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(status)); -} - EcalRechitChannelStatusGPU::Product const& EcalRechitChannelStatusGPU::getProduct(cudaStream_t cudaStream) const { auto const& product = product_.dataForCurrentDeviceAsync( cudaStream, [this](EcalRechitChannelStatusGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.status, this->status_.size() * sizeof(uint16_t))); + // allocate + product.status = cms::cuda::make_device_unique(status_.size(), cudaStream); // transfer - cudaCheck(cudaMemcpyAsync(product.status, - this->status_.data(), - this->status_.size() * sizeof(uint16_t), - cudaMemcpyHostToDevice, - cudaStream)); + cms::cuda::copyAsync(product.status, status_, cudaStream); }); return product; diff --git a/CondFormats/EcalObjects/src/EcalSamplesCorrelationGPU.cc b/CondFormats/EcalObjects/src/EcalSamplesCorrelationGPU.cc new file mode 100644 index 0000000000000..9e355045149ea --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalSamplesCorrelationGPU.cc @@ -0,0 +1,66 @@ +#include "CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalSamplesCorrelationGPU::EcalSamplesCorrelationGPU(EcalSamplesCorrelation const& values) { + EBG12SamplesCorrelation_.reserve(values.EBG12SamplesCorrelation.size()); + for (const auto& EBG12SamplesCorrelation : values.EBG12SamplesCorrelation) { + EBG12SamplesCorrelation_.emplace_back(EBG12SamplesCorrelation); + } + + EBG6SamplesCorrelation_.reserve(values.EBG6SamplesCorrelation.size()); + for (const auto& EBG6SamplesCorrelation : values.EBG6SamplesCorrelation) { + EBG6SamplesCorrelation_.emplace_back(EBG6SamplesCorrelation); + } + + EBG1SamplesCorrelation_.reserve(values.EBG1SamplesCorrelation.size()); + for (const auto& EBG1SamplesCorrelation : values.EBG1SamplesCorrelation) { + EBG1SamplesCorrelation_.emplace_back(EBG1SamplesCorrelation); + } + + EEG12SamplesCorrelation_.reserve(values.EEG12SamplesCorrelation.size()); + for (const auto& EEG12SamplesCorrelation : values.EEG12SamplesCorrelation) { + EEG12SamplesCorrelation_.emplace_back(EEG12SamplesCorrelation); + } + + EEG6SamplesCorrelation_.reserve(values.EEG6SamplesCorrelation.size()); + for (const auto& EEG6SamplesCorrelation : values.EEG6SamplesCorrelation) { + EEG6SamplesCorrelation_.emplace_back(EEG6SamplesCorrelation); + } + + EEG1SamplesCorrelation_.reserve(values.EEG1SamplesCorrelation.size()); + for (const auto& EEG1SamplesCorrelation : values.EEG1SamplesCorrelation) { + EEG1SamplesCorrelation_.emplace_back(EEG1SamplesCorrelation); + } +} + +EcalSamplesCorrelationGPU::Product const& EcalSamplesCorrelationGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalSamplesCorrelationGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.EBG12SamplesCorrelation = + cms::cuda::make_device_unique(EBG12SamplesCorrelation_.size(), cudaStream); + product.EBG6SamplesCorrelation = + cms::cuda::make_device_unique(EBG6SamplesCorrelation_.size(), cudaStream); + product.EBG1SamplesCorrelation = + cms::cuda::make_device_unique(EBG1SamplesCorrelation_.size(), cudaStream); + product.EEG12SamplesCorrelation = + cms::cuda::make_device_unique(EEG12SamplesCorrelation_.size(), cudaStream); + product.EEG6SamplesCorrelation = + cms::cuda::make_device_unique(EEG6SamplesCorrelation_.size(), cudaStream); + product.EEG1SamplesCorrelation = + cms::cuda::make_device_unique(EEG1SamplesCorrelation_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.EBG12SamplesCorrelation, EBG12SamplesCorrelation_, cudaStream); + cms::cuda::copyAsync(product.EBG6SamplesCorrelation, EBG6SamplesCorrelation_, cudaStream); + cms::cuda::copyAsync(product.EBG1SamplesCorrelation, EBG1SamplesCorrelation_, cudaStream); + cms::cuda::copyAsync(product.EEG12SamplesCorrelation, EEG12SamplesCorrelation_, cudaStream); + cms::cuda::copyAsync(product.EEG6SamplesCorrelation, EEG6SamplesCorrelation_, cudaStream); + cms::cuda::copyAsync(product.EEG1SamplesCorrelation, EEG1SamplesCorrelation_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalSamplesCorrelationGPU); diff --git a/CondFormats/EcalObjects/src/EcalTimeBiasCorrectionsGPU.cc b/CondFormats/EcalObjects/src/EcalTimeBiasCorrectionsGPU.cc new file mode 100644 index 0000000000000..268a22f7182f5 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalTimeBiasCorrectionsGPU.cc @@ -0,0 +1,53 @@ +#include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalTimeBiasCorrectionsGPU::EcalTimeBiasCorrectionsGPU(EcalTimeBiasCorrections const& values) { + ebTimeCorrAmplitudeBins_.reserve(values.EBTimeCorrAmplitudeBins.size()); + for (const auto& ebTimeCorrAmplitudeBin : values.EBTimeCorrAmplitudeBins) { + ebTimeCorrAmplitudeBins_.emplace_back(ebTimeCorrAmplitudeBin); + } + + ebTimeCorrShiftBins_.reserve(values.EBTimeCorrAmplitudeBins.size()); + for (const auto& ebTimeCorrShiftBin : values.EBTimeCorrShiftBins) { + ebTimeCorrShiftBins_.emplace_back(ebTimeCorrShiftBin); + } + + eeTimeCorrAmplitudeBins_.reserve(values.EETimeCorrAmplitudeBins.size()); + for (const auto& eeTimeCorrAmplitudeBin : values.EETimeCorrAmplitudeBins) { + eeTimeCorrAmplitudeBins_.emplace_back(eeTimeCorrAmplitudeBin); + } + + eeTimeCorrShiftBins_.reserve(values.EETimeCorrAmplitudeBins.size()); + for (const auto& eeTimeCorrShiftBin : values.EETimeCorrShiftBins) { + eeTimeCorrShiftBins_.emplace_back(eeTimeCorrShiftBin); + } +} + +EcalTimeBiasCorrectionsGPU::Product const& EcalTimeBiasCorrectionsGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalTimeBiasCorrectionsGPU::Product& product, cudaStream_t cudaStream) { + // to get the size of vectors later on + // should be removed and host conditions' objects used directly + product.ebTimeCorrAmplitudeBinsSize = this->ebTimeCorrAmplitudeBins_.size(); + product.eeTimeCorrAmplitudeBinsSize = this->eeTimeCorrAmplitudeBins_.size(); + + // allocate + product.ebTimeCorrAmplitudeBins = + cms::cuda::make_device_unique(ebTimeCorrAmplitudeBins_.size(), cudaStream); + product.ebTimeCorrShiftBins = cms::cuda::make_device_unique(ebTimeCorrShiftBins_.size(), cudaStream); + product.eeTimeCorrAmplitudeBins = + cms::cuda::make_device_unique(eeTimeCorrAmplitudeBins_.size(), cudaStream); + product.eeTimeCorrShiftBins = cms::cuda::make_device_unique(eeTimeCorrShiftBins_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.ebTimeCorrAmplitudeBins, ebTimeCorrAmplitudeBins_, cudaStream); + cms::cuda::copyAsync(product.ebTimeCorrShiftBins, ebTimeCorrShiftBins_, cudaStream); + cms::cuda::copyAsync(product.eeTimeCorrAmplitudeBins, eeTimeCorrAmplitudeBins_, cudaStream); + cms::cuda::copyAsync(product.eeTimeCorrShiftBins, eeTimeCorrShiftBins_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalTimeBiasCorrectionsGPU); diff --git a/CondFormats/EcalObjects/src/EcalTimeCalibConstantsGPU.cc b/CondFormats/EcalObjects/src/EcalTimeCalibConstantsGPU.cc new file mode 100644 index 0000000000000..80537fa9f07c8 --- /dev/null +++ b/CondFormats/EcalObjects/src/EcalTimeCalibConstantsGPU.cc @@ -0,0 +1,26 @@ +#include "CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h" + +#include "FWCore/Utilities/interface/typelookup.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" + +EcalTimeCalibConstantsGPU::EcalTimeCalibConstantsGPU(EcalTimeCalibConstants const& values) { + values_.reserve(values.size()); + for (size_t i = 0; i < values.size(); ++i) { + values_.emplace_back(values[i]); + } + offset_ = values.barrelItems().size(); +} + +EcalTimeCalibConstantsGPU::Product const& EcalTimeCalibConstantsGPU::getProduct(cudaStream_t cudaStream) const { + auto const& product = product_.dataForCurrentDeviceAsync( + cudaStream, [this](EcalTimeCalibConstantsGPU::Product& product, cudaStream_t cudaStream) { + // allocate + product.values = cms::cuda::make_device_unique(values_.size(), cudaStream); + // transfer + cms::cuda::copyAsync(product.values, values_, cudaStream); + }); + + return product; +} + +TYPELOOKUP_DATA_REG(EcalTimeCalibConstantsGPU); diff --git a/EventFilter/EcalRawToDigi/src/ElectronicsMappingGPU.cc b/CondFormats/EcalObjects/src/ElectronicsMappingGPU.cc similarity index 67% rename from EventFilter/EcalRawToDigi/src/ElectronicsMappingGPU.cc rename to CondFormats/EcalObjects/src/ElectronicsMappingGPU.cc index 8264c501a896c..343441e5d059a 100644 --- a/EventFilter/EcalRawToDigi/src/ElectronicsMappingGPU.cc +++ b/CondFormats/EcalObjects/src/ElectronicsMappingGPU.cc @@ -1,7 +1,7 @@ -#include "EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h" +#include "CondFormats/EcalObjects/interface/ElectronicsMappingGPU.h" #include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" #include "DataFormats/EcalDetId/interface/EcalElectronicsId.h" @@ -29,23 +29,13 @@ namespace ecal { } } - ElectronicsMappingGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(eid2did)); - } - ElectronicsMappingGPU::Product const& ElectronicsMappingGPU::getProduct(cudaStream_t cudaStream) const { auto const& product = product_.dataForCurrentDeviceAsync( cudaStream, [this](ElectronicsMappingGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.eid2did, this->eid2did_.size() * sizeof(uint32_t))); - + // allocate + product.eid2did = cms::cuda::make_device_unique(eid2did_.size(), cudaStream); // transfer - cudaCheck(cudaMemcpyAsync(product.eid2did, - this->eid2did_.data(), - this->eid2did_.size() * sizeof(uint32_t), - cudaMemcpyHostToDevice, - cudaStream)); + cms::cuda::copyAsync(product.eid2did, eid2did_, cudaStream); }); return product; diff --git a/EventFilter/EcalRawToDigi/BuildFile.xml b/EventFilter/EcalRawToDigi/BuildFile.xml index a00aec3dfa99b..bd9d2ea239e3a 100644 --- a/EventFilter/EcalRawToDigi/BuildFile.xml +++ b/EventFilter/EcalRawToDigi/BuildFile.xml @@ -1,5 +1,4 @@ - @@ -9,8 +8,6 @@ - - diff --git a/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h b/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h index a6429121adc82..c94cfe0c3805f 100644 --- a/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h +++ b/EventFilter/EcalRawToDigi/plugins/DeclsForKernels.h @@ -4,9 +4,9 @@ #include #include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" +#include "CondFormats/EcalObjects/interface/ElectronicsMappingGPU.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" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" diff --git a/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc b/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc index 5563dd5b52cc8..97f9828e9f5ec 100644 --- a/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc +++ b/EventFilter/EcalRawToDigi/plugins/EcalCPUDigisProducer.cc @@ -2,12 +2,12 @@ #include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" #include "CondFormats/DataRecord/interface/EcalMappingElectronicsRcd.h" +#include "CondFormats/EcalObjects/interface/ElectronicsMappingGPU.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" -#include "EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/MakerMacros.h" diff --git a/EventFilter/EcalRawToDigi/plugins/EcalRawESProducersGPUDefs.cc b/EventFilter/EcalRawToDigi/plugins/EcalRawESProducersGPUDefs.cc index 84fcc7b2b2952..b2569a57ee575 100644 --- a/EventFilter/EcalRawToDigi/plugins/EcalRawESProducersGPUDefs.cc +++ b/EventFilter/EcalRawToDigi/plugins/EcalRawESProducersGPUDefs.cc @@ -1,5 +1,5 @@ #include "CondFormats/DataRecord/interface/EcalMappingElectronicsRcd.h" -#include "EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h" +#include "CondFormats/EcalObjects/interface/ElectronicsMappingGPU.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "HeterogeneousCore/CUDACore/interface/ConvertingESProducerT.h" diff --git a/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc b/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc index 4f0743c9b1b51..5b58bf159b9d0 100644 --- a/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc +++ b/EventFilter/EcalRawToDigi/plugins/EcalRawToDigiGPU.cc @@ -1,8 +1,8 @@ #include "CUDADataFormats/EcalDigi/interface/DigisCollection.h" #include "CondFormats/DataRecord/interface/EcalMappingElectronicsRcd.h" +#include "CondFormats/EcalObjects/interface/ElectronicsMappingGPU.h" #include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" #include "DataFormats/FEDRawData/interface/FEDRawDataCollection.h" -#include "EventFilter/EcalRawToDigi/interface/ElectronicsMappingGPU.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/MakerMacros.h" diff --git a/RecoLocalCalo/EcalRecAlgos/BuildFile.xml b/RecoLocalCalo/EcalRecAlgos/BuildFile.xml index b14fd8e0787f2..723cb21a33908 100644 --- a/RecoLocalCalo/EcalRecAlgos/BuildFile.xml +++ b/RecoLocalCalo/EcalRecAlgos/BuildFile.xml @@ -1,12 +1,9 @@ - - - diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h b/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h deleted file mode 100644 index 4a6cd34fcd171..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h +++ /dev/null @@ -1,53 +0,0 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAPDPNRatiosGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAPDPNRatiosGPU_h - -#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatios.h" - -#ifndef __CUDACC__ -#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" -#include "HeterogeneousCore/CUDACore/interface/ESProduct.h" -#endif // __CUDACC__ - -class EcalLaserAPDPNRatiosGPU { -public: - struct Product { - ~Product(); - float *p1 = nullptr; - float *p2 = nullptr; - float *p3 = nullptr; - edm::TimeValue_t *t1 = nullptr; - edm::TimeValue_t *t2 = nullptr; - edm::TimeValue_t *t3 = nullptr; - }; - -#ifndef __CUDACC__ - - // - EcalLaserAPDPNRatiosGPU(EcalLaserAPDPNRatios const &); - - // will call dealloation for Product thru ~Product - ~EcalLaserAPDPNRatiosGPU() = default; - - // get device pointers - Product const &getProduct(cudaStream_t) const; - - // - static std::string name() { return std::string{"ecalLaserAPDPNRatiosGPU"}; } - -private: - // in the future, we need to arrange so to avoid this copy on the host - // store eb first then ee - std::vector > p1_; - std::vector > p2_; - std::vector > p3_; - - std::vector > t1_; - std::vector > t2_; - std::vector > t3_; - - cms::cuda::ESProduct product_; - -#endif // __CUDACC__ -}; - -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalLaserAPDPNRatiosGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h b/RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h deleted file mode 100644 index e1dee2d505e6c..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h +++ /dev/null @@ -1,44 +0,0 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalSamplesCorrelationGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalSamplesCorrelationGPU_h - -#include "CondFormats/EcalObjects/interface/EcalSamplesCorrelation.h" - -#ifndef __CUDACC__ -#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" -#include "HeterogeneousCore/CUDACore/interface/ESProduct.h" -#endif // __CUDACC__ - -class EcalSamplesCorrelationGPU { -public: - struct Product { - ~Product(); - double *EBG12SamplesCorrelation = nullptr, *EBG6SamplesCorrelation = nullptr, *EBG1SamplesCorrelation = nullptr; - double *EEG12SamplesCorrelation = nullptr, *EEG6SamplesCorrelation = nullptr, *EEG1SamplesCorrelation = nullptr; - }; - -#ifndef __CUDACC__ - // rearrange pedestals - EcalSamplesCorrelationGPU(EcalSamplesCorrelation const&); - - // will call dealloation for Product thru ~Product - ~EcalSamplesCorrelationGPU() = default; - - // get device pointers - Product const& getProduct(cudaStream_t) const; - - // - static std::string name() { return std::string{"ecalSamplesCorrelationGPU"}; } - -private: - std::vector const& EBG12SamplesCorrelation_; - std::vector const& EBG6SamplesCorrelation_; - std::vector const& EBG1SamplesCorrelation_; - std::vector const& EEG12SamplesCorrelation_; - std::vector const& EEG6SamplesCorrelation_; - std::vector const& EEG1SamplesCorrelation_; - - cms::cuda::ESProduct product_; -#endif // __CUDACC__ -}; - -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalSamplesCorrelationGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h b/RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h deleted file mode 100644 index 9e2bf0aa18909..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h +++ /dev/null @@ -1,49 +0,0 @@ -#ifndef RecoLocalCalo_EcalRecAlgos_interface_EcalTimeBiasCorrectionsGPU_h -#define RecoLocalCalo_EcalRecAlgos_interface_EcalTimeBiasCorrectionsGPU_h - -#include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrections.h" - -#ifndef __CUDACC__ -#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h" -#include "HeterogeneousCore/CUDACore/interface/ESProduct.h" -#endif // __CUDACC__ - -class EcalTimeBiasCorrectionsGPU { -public: - struct Product { - ~Product(); - float *EBTimeCorrAmplitudeBins, *EBTimeCorrShiftBins; - float *EETimeCorrAmplitudeBins, *EETimeCorrShiftBins; - int EBTimeCorrAmplitudeBinsSize, EETimeCorrAmplitudeBinsSize; - }; - - // rearrange pedestals - EcalTimeBiasCorrectionsGPU(EcalTimeBiasCorrections const&); - -#ifndef __CUDACC__ - - // will call dealloation for Product thru ~Product - ~EcalTimeBiasCorrectionsGPU() = default; - - // get device pointers - Product const& getProduct(cudaStream_t) const; - - // - static std::string name() { return std::string{"ecalTimeBiasCorrectionsGPU"}; } -#endif // __CUDACC__ - - std::vector const& EBTimeCorrAmplitudeBins() const { return EBTimeCorrAmplitudeBins_; } - std::vector const& EETimeCorrAmplitudeBins() const { return EETimeCorrAmplitudeBins_; } - -private: - std::vector const& EBTimeCorrAmplitudeBins_; - std::vector const& EBTimeCorrShiftBins_; - std::vector const& EETimeCorrAmplitudeBins_; - std::vector const& EETimeCorrShiftBins_; - -#ifndef __CUDACC__ - cms::cuda::ESProduct product_; -#endif // __CUDACC__ -}; - -#endif // RecoLocalCalo_EcalRecAlgos_interface_EcalTimeBiasCorrectionsGPU_h diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalGainRatiosGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalGainRatiosGPU.cc deleted file mode 100644 index d5980d8a757aa..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalGainRatiosGPU.cc +++ /dev/null @@ -1,52 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalGainRatiosGPU::EcalGainRatiosGPU(EcalGainRatios const& values) - : gain12Over6_(values.size()), gain6Over1_(values.size()) { - // fill in eb - auto const& barrelValues = values.barrelItems(); - for (unsigned int i = 0; i < barrelValues.size(); i++) { - gain12Over6_[i] = barrelValues[i].gain12Over6(); - gain6Over1_[i] = barrelValues[i].gain6Over1(); - } - - // fill in ee - auto const& endcapValues = values.endcapItems(); - auto const offset = barrelValues.size(); - for (unsigned int i = 0; i < endcapValues.size(); i++) { - gain12Over6_[offset + i] = endcapValues[i].gain12Over6(); - gain6Over1_[offset + i] = endcapValues[i].gain6Over1(); - } -} - -EcalGainRatiosGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(gain12Over6)); - cudaCheck(cudaFree(gain6Over1)); -} - -EcalGainRatiosGPU::Product const& EcalGainRatiosGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalGainRatiosGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.gain12Over6, this->gain12Over6_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.gain6Over1, this->gain6Over1_.size() * sizeof(float))); - // transfer - cudaCheck(cudaMemcpyAsync(product.gain12Over6, - this->gain12Over6_.data(), - this->gain12Over6_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.gain6Over1, - this->gain6Over1_.data(), - this->gain6Over1_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalGainRatiosGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalIntercalibConstantsGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalIntercalibConstantsGPU.cc deleted file mode 100644 index dec10cff57dd0..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalIntercalibConstantsGPU.cc +++ /dev/null @@ -1,40 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalIntercalibConstantsGPU::EcalIntercalibConstantsGPU(EcalIntercalibConstants const& values) - : valuesEB_{values.barrelItems()}, valuesEE_{values.endcapItems()} {} - -EcalIntercalibConstantsGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(values)); -} - -EcalIntercalibConstantsGPU::Product const& EcalIntercalibConstantsGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalIntercalibConstantsGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck( - cudaMalloc((void**)&product.values, (this->valuesEB_.size() + this->valuesEE_.size()) * sizeof(float))); - - // offset in floats, not bytes - auto const offset = this->valuesEB_.size(); - - // transfer - cudaCheck(cudaMemcpyAsync(product.values, - this->valuesEB_.data(), - this->valuesEB_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.values + offset, - this->valuesEE_.data(), - this->valuesEE_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalIntercalibConstantsGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAPDPNRatiosGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAPDPNRatiosGPU.cc deleted file mode 100644 index 4aa92ea6750fe..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAPDPNRatiosGPU.cc +++ /dev/null @@ -1,86 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalLaserAPDPNRatiosGPU::EcalLaserAPDPNRatiosGPU(EcalLaserAPDPNRatios const& values) - : p1_(values.getLaserMap().size()), - p2_(values.getLaserMap().size()), - p3_(values.getLaserMap().size()), - t1_(values.getTimeMap().size()), - t2_(values.getTimeMap().size()), - t3_(values.getTimeMap().size()) { - // fill in eb - // auto const& barrelValues = values.barrelItems(); - for (unsigned int i = 0; i < values.getLaserMap().barrelItems().size(); i++) { - p1_[i] = values.getLaserMap().barrelItems()[i].p1; - p2_[i] = values.getLaserMap().barrelItems()[i].p2; - p3_[i] = values.getLaserMap().barrelItems()[i].p3; - } - - // fill in ee - // auto const& endcapValues = values.endcapItems(); - auto const offset_laser = values.getLaserMap().barrelItems().size(); - for (unsigned int i = 0; i < values.getLaserMap().endcapItems().size(); i++) { - p1_[offset_laser + i] = values.getLaserMap().endcapItems()[i].p1; - p2_[offset_laser + i] = values.getLaserMap().endcapItems()[i].p2; - p3_[offset_laser + i] = values.getLaserMap().endcapItems()[i].p3; - } - - // Time is a simple std::vector - // typedef std::vector EcalLaserTimeStampMap; - for (unsigned int i = 0; i < values.getTimeMap().size(); i++) { - t1_[i] = values.getTimeMap()[i].t1.value(); - t2_[i] = values.getTimeMap()[i].t2.value(); - t3_[i] = values.getTimeMap()[i].t3.value(); - } -} - -EcalLaserAPDPNRatiosGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(p1)); - cudaCheck(cudaFree(p2)); - cudaCheck(cudaFree(p3)); - cudaCheck(cudaFree(t1)); - cudaCheck(cudaFree(t2)); - cudaCheck(cudaFree(t3)); -} - -EcalLaserAPDPNRatiosGPU::Product const& EcalLaserAPDPNRatiosGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalLaserAPDPNRatiosGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.p1, this->p1_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.p2, this->p2_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.p3, this->p3_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.t1, this->t1_.size() * sizeof(edm::TimeValue_t))); - cudaCheck(cudaMalloc((void**)&product.t2, this->t2_.size() * sizeof(edm::TimeValue_t))); - cudaCheck(cudaMalloc((void**)&product.t3, this->t3_.size() * sizeof(edm::TimeValue_t))); - // transfer - cudaCheck(cudaMemcpyAsync( - product.p1, this->p1_.data(), this->p1_.size() * sizeof(float), cudaMemcpyHostToDevice, cudaStream)); - cudaCheck(cudaMemcpyAsync( - product.p2, this->p2_.data(), this->p2_.size() * sizeof(float), cudaMemcpyHostToDevice, cudaStream)); - cudaCheck(cudaMemcpyAsync( - product.p3, this->p3_.data(), this->p3_.size() * sizeof(float), cudaMemcpyHostToDevice, cudaStream)); - cudaCheck(cudaMemcpyAsync(product.t1, - this->t1_.data(), - this->t1_.size() * sizeof(edm::TimeValue_t), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.t2, - this->t2_.data(), - this->t2_.size() * sizeof(edm::TimeValue_t), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.t3, - this->t3_.data(), - this->t3_.size() * sizeof(edm::TimeValue_t), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalLaserAPDPNRatiosGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAPDPNRatiosRefGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAPDPNRatiosRefGPU.cc deleted file mode 100644 index 8f77cf48fe1d1..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAPDPNRatiosRefGPU.cc +++ /dev/null @@ -1,40 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalLaserAPDPNRatiosRefGPU::EcalLaserAPDPNRatiosRefGPU(EcalLaserAPDPNRatiosRef const& values) - : valuesEB_{values.barrelItems()}, valuesEE_{values.endcapItems()} {} - -EcalLaserAPDPNRatiosRefGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(values)); -} - -EcalLaserAPDPNRatiosRefGPU::Product const& EcalLaserAPDPNRatiosRefGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalLaserAPDPNRatiosRefGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck( - cudaMalloc((void**)&product.values, (this->valuesEB_.size() + this->valuesEE_.size()) * sizeof(float))); - - // offset in floats, not bytes - auto const offset = this->valuesEB_.size(); - - // transfer - cudaCheck(cudaMemcpyAsync(product.values, - this->valuesEB_.data(), - this->valuesEB_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.values + offset, - this->valuesEE_.data(), - this->valuesEE_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalLaserAPDPNRatiosRefGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAlphasGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAlphasGPU.cc deleted file mode 100644 index 91de441bff683..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalLaserAlphasGPU.cc +++ /dev/null @@ -1,40 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalLaserAlphasGPU::EcalLaserAlphasGPU(EcalLaserAlphas const& values) - : valuesEB_{values.barrelItems()}, valuesEE_{values.endcapItems()} {} - -EcalLaserAlphasGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(values)); -} - -EcalLaserAlphasGPU::Product const& EcalLaserAlphasGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalLaserAlphasGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck( - cudaMalloc((void**)&product.values, (this->valuesEB_.size() + this->valuesEE_.size()) * sizeof(float))); - - // offset in floats, not bytes - auto const offset = this->valuesEB_.size(); - - // transfer - cudaCheck(cudaMemcpyAsync(product.values, - this->valuesEB_.data(), - this->valuesEB_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.values + offset, - this->valuesEE_.data(), - this->valuesEE_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalLaserAlphasGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalLinearCorrectionsGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalLinearCorrectionsGPU.cc deleted file mode 100644 index 0af2a9044ab65..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalLinearCorrectionsGPU.cc +++ /dev/null @@ -1,84 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalLinearCorrectionsGPU::EcalLinearCorrectionsGPU(EcalLinearCorrections const& values) - : p1_(values.getValueMap().size()), - p2_(values.getValueMap().size()), - p3_(values.getValueMap().size()), - t1_(values.getTimeMap().size()), - t2_(values.getTimeMap().size()), - t3_(values.getTimeMap().size()) { - // fill in eb - for (unsigned int i = 0; i < values.getValueMap().barrelItems().size(); i++) { - p1_[i] = values.getValueMap().barrelItems()[i].p1; - p2_[i] = values.getValueMap().barrelItems()[i].p2; - p3_[i] = values.getValueMap().barrelItems()[i].p3; - } - - // fill in ee - auto const offset_laser = values.getValueMap().barrelItems().size(); - for (unsigned int i = 0; i < values.getValueMap().endcapItems().size(); i++) { - p1_[offset_laser + i] = values.getValueMap().endcapItems()[i].p1; - p2_[offset_laser + i] = values.getValueMap().endcapItems()[i].p2; - p3_[offset_laser + i] = values.getValueMap().endcapItems()[i].p3; - } - - // Time is a simple std::vector - // typedef std::vector EcalLaserTimeStampMap; - for (unsigned int i = 0; i < values.getTimeMap().size(); i++) { - t1_[i] = values.getTimeMap()[i].t1.value(); - t2_[i] = values.getTimeMap()[i].t2.value(); - t3_[i] = values.getTimeMap()[i].t3.value(); - } -} - -EcalLinearCorrectionsGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(p1)); - cudaCheck(cudaFree(p2)); - cudaCheck(cudaFree(p3)); - cudaCheck(cudaFree(t1)); - cudaCheck(cudaFree(t2)); - cudaCheck(cudaFree(t3)); -} - -EcalLinearCorrectionsGPU::Product const& EcalLinearCorrectionsGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalLinearCorrectionsGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.p1, this->p1_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.p2, this->p2_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.p3, this->p3_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.t1, this->t1_.size() * sizeof(edm::TimeValue_t))); - cudaCheck(cudaMalloc((void**)&product.t2, this->t2_.size() * sizeof(edm::TimeValue_t))); - cudaCheck(cudaMalloc((void**)&product.t3, this->t3_.size() * sizeof(edm::TimeValue_t))); - // transfer - cudaCheck(cudaMemcpyAsync( - product.p1, this->p1_.data(), this->p1_.size() * sizeof(float), cudaMemcpyHostToDevice, cudaStream)); - cudaCheck(cudaMemcpyAsync( - product.p2, this->p2_.data(), this->p2_.size() * sizeof(float), cudaMemcpyHostToDevice, cudaStream)); - cudaCheck(cudaMemcpyAsync( - product.p3, this->p3_.data(), this->p3_.size() * sizeof(float), cudaMemcpyHostToDevice, cudaStream)); - cudaCheck(cudaMemcpyAsync(product.t1, - this->t1_.data(), - this->t1_.size() * sizeof(edm::TimeValue_t), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.t2, - this->t2_.data(), - this->t2_.size() * sizeof(edm::TimeValue_t), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.t3, - this->t3_.data(), - this->t3_.size() * sizeof(edm::TimeValue_t), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalLinearCorrectionsGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalMultifitParametersGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalMultifitParametersGPU.cc deleted file mode 100644 index 010da6444b614..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalMultifitParametersGPU.cc +++ /dev/null @@ -1,66 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalMultifitParametersGPU::EcalMultifitParametersGPU(edm::ParameterSet const& ps) { - auto const& amplitudeFitParametersEB = ps.getParameter>("EBamplitudeFitParameters"); - auto const& amplitudeFitParametersEE = ps.getParameter>("EEamplitudeFitParameters"); - auto const& timeFitParametersEB = ps.getParameter>("EBtimeFitParameters"); - auto const& timeFitParametersEE = ps.getParameter>("EEtimeFitParameters"); - - amplitudeFitParametersEB_.resize(amplitudeFitParametersEB.size()); - amplitudeFitParametersEE_.resize(amplitudeFitParametersEE.size()); - timeFitParametersEB_.resize(timeFitParametersEB.size()); - timeFitParametersEE_.resize(timeFitParametersEE.size()); - - std::copy(amplitudeFitParametersEB.begin(), amplitudeFitParametersEB.end(), amplitudeFitParametersEB_.begin()); - std::copy(amplitudeFitParametersEE.begin(), amplitudeFitParametersEE.end(), amplitudeFitParametersEE_.begin()); - std::copy(timeFitParametersEB.begin(), timeFitParametersEB.end(), timeFitParametersEB_.begin()); - std::copy(timeFitParametersEE.begin(), timeFitParametersEE.end(), timeFitParametersEE_.begin()); -} - -EcalMultifitParametersGPU::Product::~Product() { - cudaCheck(cudaFree(amplitudeFitParametersEB)); - cudaCheck(cudaFree(amplitudeFitParametersEE)); - cudaCheck(cudaFree(timeFitParametersEB)); - cudaCheck(cudaFree(timeFitParametersEE)); -} - -EcalMultifitParametersGPU::Product const& EcalMultifitParametersGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalMultifitParametersGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.amplitudeFitParametersEB, - this->amplitudeFitParametersEB_.size() * sizeof(double))); - cudaCheck(cudaMalloc((void**)&product.amplitudeFitParametersEE, - this->amplitudeFitParametersEE_.size() * sizeof(double))); - cudaCheck(cudaMalloc((void**)&product.timeFitParametersEB, this->timeFitParametersEB_.size() * sizeof(double))); - cudaCheck(cudaMalloc((void**)&product.timeFitParametersEE, this->timeFitParametersEE_.size() * sizeof(double))); - - // transfer - cudaCheck(cudaMemcpyAsync(product.amplitudeFitParametersEB, - this->amplitudeFitParametersEB_.data(), - this->amplitudeFitParametersEB_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.amplitudeFitParametersEE, - this->amplitudeFitParametersEE_.data(), - this->amplitudeFitParametersEE_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.timeFitParametersEB, - this->timeFitParametersEB_.data(), - this->timeFitParametersEB_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.timeFitParametersEE, - this->timeFitParametersEE_.data(), - this->timeFitParametersEE_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - }); - return product; -} - -TYPELOOKUP_DATA_REG(EcalMultifitParametersGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalPedestalsGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalPedestalsGPU.cc deleted file mode 100644 index 9e3284cd9c7c8..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalPedestalsGPU.cc +++ /dev/null @@ -1,94 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalPedestalsGPU::EcalPedestalsGPU(EcalPedestals const& pedestals) - : mean_x12_(pedestals.size()), - rms_x12_(pedestals.size()), - mean_x6_(pedestals.size()), - rms_x6_(pedestals.size()), - mean_x1_(pedestals.size()), - rms_x1_(pedestals.size()) { - // fill in eb - auto const& barrelValues = pedestals.barrelItems(); - for (unsigned int i = 0; i < barrelValues.size(); i++) { - mean_x12_[i] = barrelValues[i].mean_x12; - rms_x12_[i] = barrelValues[i].rms_x12; - mean_x6_[i] = barrelValues[i].mean_x6; - rms_x6_[i] = barrelValues[i].rms_x6; - mean_x1_[i] = barrelValues[i].mean_x1; - rms_x1_[i] = barrelValues[i].rms_x1; - } - - // fill in ee - auto const& endcapValues = pedestals.endcapItems(); - auto const offset = barrelValues.size(); - for (unsigned int i = 0; i < endcapValues.size(); i++) { - mean_x12_[offset + i] = endcapValues[i].mean_x12; - rms_x12_[offset + i] = endcapValues[i].rms_x12; - mean_x6_[offset + i] = endcapValues[i].mean_x6; - rms_x6_[offset + i] = endcapValues[i].rms_x6; - mean_x1_[offset + i] = endcapValues[i].mean_x1; - rms_x1_[offset + i] = endcapValues[i].rms_x1; - } -} - -EcalPedestalsGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(mean_x12)); - cudaCheck(cudaFree(rms_x12)); - cudaCheck(cudaFree(mean_x6)); - cudaCheck(cudaFree(rms_x6)); - cudaCheck(cudaFree(mean_x1)); - cudaCheck(cudaFree(rms_x1)); -} - -EcalPedestalsGPU::Product const& EcalPedestalsGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalPedestalsGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.mean_x12, this->mean_x12_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.rms_x12, this->mean_x12_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.mean_x6, this->mean_x12_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.rms_x6, this->mean_x12_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.mean_x1, this->mean_x12_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.rms_x1, this->mean_x12_.size() * sizeof(float))); - - // transfer - cudaCheck(cudaMemcpyAsync(product.mean_x12, - this->mean_x12_.data(), - this->mean_x12_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.rms_x12, - this->rms_x12_.data(), - this->rms_x12_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.mean_x6, - this->mean_x6_.data(), - this->mean_x6_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.rms_x6, - this->rms_x6_.data(), - this->rms_x6_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.mean_x1, - this->mean_x1_.data(), - this->mean_x1_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.rms_x1, - this->rms_x1_.data(), - this->rms_x1_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalPedestalsGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalRecHitParametersGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalRecHitParametersGPU.cc deleted file mode 100644 index 0f6812d6d6ffe..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalRecHitParametersGPU.cc +++ /dev/null @@ -1,82 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRecHitParametersGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "CommonTools/Utils/interface/StringToEnumValue.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h" -#include "DataFormats/EcalRecHit/interface/EcalRecHit.h" - -EcalRecHitParametersGPU::EcalRecHitParametersGPU(edm::ParameterSet const& ps) { - auto const& ChannelStatusToBeExcluded = StringToEnumValue( - ps.getParameter>("ChannelStatusToBeExcluded")); - - ChannelStatusToBeExcluded_.resize(ChannelStatusToBeExcluded.size()); - std::copy(ChannelStatusToBeExcluded.begin(), ChannelStatusToBeExcluded.end(), ChannelStatusToBeExcluded_.begin()); - - // https://github.com/cms-sw/cmssw/blob/266e21cfc9eb409b093e4cf064f4c0a24c6ac293/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitWorkerSimple.cc - - // Traslate string representation of flagsMapDBReco into enum values - const edm::ParameterSet& p = ps.getParameter("flagsMapDBReco"); - std::vector recoflagbitsStrings = p.getParameterNames(); - - for (unsigned int i = 0; i != recoflagbitsStrings.size(); ++i) { - EcalRecHit::Flags recoflagbit = (EcalRecHit::Flags)StringToEnumValue(recoflagbitsStrings[i]); - std::vector dbstatus_s = p.getParameter>(recoflagbitsStrings[i]); - // std::vector dbstatuses; - for (unsigned int j = 0; j != dbstatus_s.size(); ++j) { - EcalChannelStatusCode::Code dbstatus = - (EcalChannelStatusCode::Code)StringToEnumValue(dbstatus_s[j]); - expanded_v_DB_reco_flags_.push_back(dbstatus); - } - - expanded_Sizes_v_DB_reco_flags_.push_back(dbstatus_s.size()); - expanded_flagbit_v_DB_reco_flags_.push_back(recoflagbit); - } -} - -EcalRecHitParametersGPU::Product::~Product() { - cudaCheck(cudaFree(ChannelStatusToBeExcluded)); - cudaCheck(cudaFree(expanded_v_DB_reco_flags)); - cudaCheck(cudaFree(expanded_Sizes_v_DB_reco_flags)); - cudaCheck(cudaFree(expanded_flagbit_v_DB_reco_flags)); -} - -EcalRecHitParametersGPU::Product const& EcalRecHitParametersGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalRecHitParametersGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.ChannelStatusToBeExcluded, - this->ChannelStatusToBeExcluded_.size() * sizeof(int))); - cudaCheck(cudaMalloc((void**)&product.expanded_v_DB_reco_flags, - this->expanded_v_DB_reco_flags_.size() * sizeof(int))); - cudaCheck(cudaMalloc((void**)&product.expanded_Sizes_v_DB_reco_flags, - this->expanded_Sizes_v_DB_reco_flags_.size() * sizeof(uint32_t))); - cudaCheck(cudaMalloc((void**)&product.expanded_flagbit_v_DB_reco_flags, - this->expanded_flagbit_v_DB_reco_flags_.size() * sizeof(uint32_t))); - - // transfer - cudaCheck(cudaMemcpyAsync(product.ChannelStatusToBeExcluded, - this->ChannelStatusToBeExcluded_.data(), - this->ChannelStatusToBeExcluded_.size() * sizeof(int), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.expanded_v_DB_reco_flags, - this->expanded_v_DB_reco_flags_.data(), - this->expanded_v_DB_reco_flags_.size() * sizeof(int), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.expanded_Sizes_v_DB_reco_flags, - this->expanded_Sizes_v_DB_reco_flags_.data(), - this->expanded_Sizes_v_DB_reco_flags_.size() * sizeof(uint32_t), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.expanded_flagbit_v_DB_reco_flags, - this->expanded_flagbit_v_DB_reco_flags_.data(), - this->expanded_flagbit_v_DB_reco_flags_.size() * sizeof(uint32_t), - cudaMemcpyHostToDevice, - cudaStream)); - }); - return product; -} - -TYPELOOKUP_DATA_REG(EcalRecHitParametersGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalSamplesCorrelationGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalSamplesCorrelationGPU.cc deleted file mode 100644 index 2a98067f51d9e..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalSamplesCorrelationGPU.cc +++ /dev/null @@ -1,76 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalSamplesCorrelationGPU::EcalSamplesCorrelationGPU(EcalSamplesCorrelation const& values) - : EBG12SamplesCorrelation_{values.EBG12SamplesCorrelation}, - EBG6SamplesCorrelation_{values.EBG6SamplesCorrelation}, - EBG1SamplesCorrelation_{values.EBG1SamplesCorrelation}, - EEG12SamplesCorrelation_{values.EEG12SamplesCorrelation}, - EEG6SamplesCorrelation_{values.EEG6SamplesCorrelation}, - EEG1SamplesCorrelation_{values.EEG1SamplesCorrelation} {} - -EcalSamplesCorrelationGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(EBG12SamplesCorrelation)); - cudaCheck(cudaFree(EBG6SamplesCorrelation)); - cudaCheck(cudaFree(EBG1SamplesCorrelation)); - cudaCheck(cudaFree(EEG12SamplesCorrelation)); - cudaCheck(cudaFree(EEG6SamplesCorrelation)); - cudaCheck(cudaFree(EEG1SamplesCorrelation)); -} - -EcalSamplesCorrelationGPU::Product const& EcalSamplesCorrelationGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalSamplesCorrelationGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck(cudaMalloc((void**)&product.EBG12SamplesCorrelation, - this->EBG12SamplesCorrelation_.size() * sizeof(double))); - cudaCheck( - cudaMalloc((void**)&product.EBG6SamplesCorrelation, this->EBG6SamplesCorrelation_.size() * sizeof(double))); - cudaCheck( - cudaMalloc((void**)&product.EBG1SamplesCorrelation, this->EBG1SamplesCorrelation_.size() * sizeof(double))); - cudaCheck(cudaMalloc((void**)&product.EEG12SamplesCorrelation, - this->EEG12SamplesCorrelation_.size() * sizeof(double))); - cudaCheck( - cudaMalloc((void**)&product.EEG6SamplesCorrelation, this->EEG6SamplesCorrelation_.size() * sizeof(double))); - cudaCheck( - cudaMalloc((void**)&product.EEG1SamplesCorrelation, this->EEG1SamplesCorrelation_.size() * sizeof(double))); - // transfer - cudaCheck(cudaMemcpyAsync(product.EBG12SamplesCorrelation, - this->EBG12SamplesCorrelation_.data(), - this->EBG12SamplesCorrelation_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EBG6SamplesCorrelation, - this->EBG6SamplesCorrelation_.data(), - this->EBG6SamplesCorrelation_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EBG1SamplesCorrelation, - this->EBG1SamplesCorrelation_.data(), - this->EBG1SamplesCorrelation_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EEG12SamplesCorrelation, - this->EEG12SamplesCorrelation_.data(), - this->EEG12SamplesCorrelation_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EEG6SamplesCorrelation, - this->EEG6SamplesCorrelation_.data(), - this->EEG6SamplesCorrelation_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EEG1SamplesCorrelation, - this->EEG1SamplesCorrelation_.data(), - this->EEG1SamplesCorrelation_.size() * sizeof(double), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalSamplesCorrelationGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalTimeBiasCorrectionsGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalTimeBiasCorrectionsGPU.cc deleted file mode 100644 index 9ab0a6302a9c4..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalTimeBiasCorrectionsGPU.cc +++ /dev/null @@ -1,61 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalTimeBiasCorrectionsGPU::EcalTimeBiasCorrectionsGPU(EcalTimeBiasCorrections const& values) - : EBTimeCorrAmplitudeBins_{values.EBTimeCorrAmplitudeBins}, - EBTimeCorrShiftBins_{values.EBTimeCorrShiftBins}, - EETimeCorrAmplitudeBins_{values.EETimeCorrAmplitudeBins}, - EETimeCorrShiftBins_{values.EETimeCorrShiftBins} {} - -EcalTimeBiasCorrectionsGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(EBTimeCorrAmplitudeBins)); - cudaCheck(cudaFree(EBTimeCorrShiftBins)); - cudaCheck(cudaFree(EETimeCorrAmplitudeBins)); - cudaCheck(cudaFree(EETimeCorrShiftBins)); -} - -EcalTimeBiasCorrectionsGPU::Product const& EcalTimeBiasCorrectionsGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalTimeBiasCorrectionsGPU::Product& product, cudaStream_t cudaStream) { - // to get the size of vectors later on - // should be removed and host conditions' objects used directly - product.EBTimeCorrAmplitudeBinsSize = this->EBTimeCorrAmplitudeBins_.size(); - product.EETimeCorrAmplitudeBinsSize = this->EETimeCorrAmplitudeBins_.size(); - - // malloc - cudaCheck(cudaMalloc((void**)&product.EBTimeCorrAmplitudeBins, - this->EBTimeCorrAmplitudeBins_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.EBTimeCorrShiftBins, this->EBTimeCorrShiftBins_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.EETimeCorrAmplitudeBins, - this->EETimeCorrAmplitudeBins_.size() * sizeof(float))); - cudaCheck(cudaMalloc((void**)&product.EETimeCorrShiftBins, this->EETimeCorrShiftBins_.size() * sizeof(float))); - // transfer - cudaCheck(cudaMemcpyAsync(product.EBTimeCorrAmplitudeBins, - this->EBTimeCorrAmplitudeBins_.data(), - this->EBTimeCorrAmplitudeBins_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EBTimeCorrShiftBins, - this->EBTimeCorrShiftBins_.data(), - this->EBTimeCorrShiftBins_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EETimeCorrAmplitudeBins, - this->EETimeCorrAmplitudeBins_.data(), - this->EETimeCorrAmplitudeBins_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.EETimeCorrShiftBins, - this->EETimeCorrShiftBins_.data(), - this->EETimeCorrShiftBins_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalTimeBiasCorrectionsGPU); diff --git a/RecoLocalCalo/EcalRecAlgos/src/EcalTimeCalibConstantsGPU.cc b/RecoLocalCalo/EcalRecAlgos/src/EcalTimeCalibConstantsGPU.cc deleted file mode 100644 index d724a33f1d4e1..0000000000000 --- a/RecoLocalCalo/EcalRecAlgos/src/EcalTimeCalibConstantsGPU.cc +++ /dev/null @@ -1,40 +0,0 @@ -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h" - -#include "FWCore/Utilities/interface/typelookup.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" - -EcalTimeCalibConstantsGPU::EcalTimeCalibConstantsGPU(EcalTimeCalibConstants const& values) - : valuesEB_{values.barrelItems()}, valuesEE_{values.endcapItems()} {} - -EcalTimeCalibConstantsGPU::Product::~Product() { - // deallocation - cudaCheck(cudaFree(values)); -} - -EcalTimeCalibConstantsGPU::Product const& EcalTimeCalibConstantsGPU::getProduct(cudaStream_t cudaStream) const { - auto const& product = product_.dataForCurrentDeviceAsync( - cudaStream, [this](EcalTimeCalibConstantsGPU::Product& product, cudaStream_t cudaStream) { - // malloc - cudaCheck( - cudaMalloc((void**)&product.values, (this->valuesEB_.size() + this->valuesEE_.size()) * sizeof(float))); - - // offset in floats, not bytes - auto const offset = this->valuesEB_.size(); - - // transfer - cudaCheck(cudaMemcpyAsync(product.values, - this->valuesEB_.data(), - this->valuesEB_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - cudaCheck(cudaMemcpyAsync(product.values + offset, - this->valuesEE_.data(), - this->valuesEE_.size() * sizeof(float), - cudaMemcpyHostToDevice, - cudaStream)); - }); - - return product; -} - -TYPELOOKUP_DATA_REG(EcalTimeCalibConstantsGPU); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.cu index f1b1a53a78a30..5a6bb0577c45c 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.cu @@ -7,6 +7,7 @@ #include "CondFormats/EcalObjects/interface/EcalPulseShapes.h" #include "CondFormats/EcalObjects/interface/EcalSamplesCorrelation.h" #include "DataFormats/EcalDigi/interface/EcalDataFrame.h" +#include "DataFormats/EcalDigi/interface/EcalMGPASample.h" #include "DataFormats/EcalRecHit/interface/EcalUncalibratedRecHit.h" #include "DataFormats/Math/interface/approx_exp.h" #include "DataFormats/Math/interface/approx_log.h" @@ -113,8 +114,8 @@ namespace ecal { // // amplitudes // - int const adc = ecal::mgpa::adc(digis_in[inputTx]); - int const gainId = ecal::mgpa::gainId(digis_in[inputTx]); + int const adc = ecalMGPA::adc(digis_in[inputTx]); + int const gainId = ecalMGPA::gainId(digis_in[inputTx]); SampleVector::Scalar amplitude = 0.; SampleVector::Scalar pedestal = 0.; SampleVector::Scalar gainratio = 0.; diff --git a/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.h index 479c623e83f62..1797fb6d2ec88 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationCommonKernels.h @@ -1,7 +1,6 @@ #ifndef RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationCommonKernels_h #define RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationCommonKernels_h -#include "Common.h" #include "DeclsForKernels.h" #include "EigenMatrixTypes_gpu.h" diff --git a/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationKernels.h index b8202f75b653b..72ccf3b11a987 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/AmplitudeComputationKernels.h @@ -1,7 +1,6 @@ #ifndef RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationKernels_h #define RecoLocalCalo_EcalRecProducers_plugins_AmplitudeComputationKernels_h -#include "Common.h" #include "DeclsForKernels.h" #include "EigenMatrixTypes_gpu.h" diff --git a/RecoLocalCalo/EcalRecProducers/plugins/Common.h b/RecoLocalCalo/EcalRecProducers/plugins/Common.h deleted file mode 100644 index 55f5f613ed356..0000000000000 --- a/RecoLocalCalo/EcalRecProducers/plugins/Common.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef RecoLocalCalo_EcalRecProducers_plugins_Common_h -#define RecoLocalCalo_EcalRecProducers_plugins_Common_h - -// a workaround for std::abs not being a constexpr function -namespace ecal { - - // temporary - namespace mgpa { - - constexpr int adc(uint16_t sample) { return sample & 0xfff; } - constexpr int gainId(uint16_t sample) { return (sample >> 12) & 0x3; } - - } // namespace mgpa - -} // namespace ecal - -#endif // RecoLocalCalo_EcalRecProducers_plugins_Common_h diff --git a/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h index cac63b6b30112..b1e1dafdb7496 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/DeclsForKernels.h @@ -13,26 +13,26 @@ #include "CondFormats/EcalObjects/interface/EcalChannelStatus.h" #include "CondFormats/EcalObjects/interface/EcalChannelStatusCode.h" #include "CondFormats/EcalObjects/interface/EcalGainRatios.h" +#include "CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h" +#include "CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h" #include "CondFormats/EcalObjects/interface/EcalPedestals.h" +#include "CondFormats/EcalObjects/interface/EcalPedestalsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h" +#include "CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h" #include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrections.h" +#include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h" #include "CondFormats/EcalObjects/interface/EcalTimeOffsetConstant.h" #include "CondFormats/EcalObjects/interface/EcalWeightSet.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h" #include "EigenMatrixTypes_gpu.h" @@ -67,11 +67,11 @@ namespace ecal { struct ConfigurationParameters { using type = double; // device ptrs - type *amplitudeFitParametersEB = nullptr, *amplitudeFitParametersEE = nullptr; + const type *amplitudeFitParametersEB = nullptr, *amplitudeFitParametersEE = nullptr; uint32_t timeFitParametersSizeEB, timeFitParametersSizeEE; // device ptrs - type *timeFitParametersEB = nullptr, *timeFitParametersEE = nullptr; + const type *timeFitParametersEB = nullptr, *timeFitParametersEE = nullptr; type timeFitLimitsFirstEB, timeFitLimitsFirstEE; type timeFitLimitsSecondEB, timeFitLimitsSecondEE; @@ -250,7 +250,7 @@ namespace ecal { // parameters that are read in the configuration file for rechit producer struct ConfigurationParameters { // device ptrs - int* ChannelStatusToBeExcluded = nullptr; + const int* ChannelStatusToBeExcluded = nullptr; uint32_t ChannelStatusToBeExcludedSize; bool killDeadChannels; @@ -267,9 +267,9 @@ namespace ecal { float EBLaserMAX; float EELaserMAX; - int* expanded_v_DB_reco_flags; - uint32_t* expanded_Sizes_v_DB_reco_flags; - uint32_t* expanded_flagbit_v_DB_reco_flags; + const int* expanded_v_DB_reco_flags; + const uint32_t* expanded_Sizes_v_DB_reco_flags; + const uint32_t* expanded_flagbit_v_DB_reco_flags; uint32_t expanded_v_DB_reco_flagsSize; uint32_t flagmask; diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalESProducersGPUDefs.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalESProducersGPUDefs.cc index 3118d54c6a7e9..5bdadecc75ef8 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalESProducersGPUDefs.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalESProducersGPUDefs.cc @@ -12,22 +12,22 @@ #include "CondFormats/DataRecord/interface/EcalSamplesCorrelationRcd.h" #include "CondFormats/DataRecord/interface/EcalTimeBiasCorrectionsRcd.h" #include "CondFormats/DataRecord/interface/EcalTimeCalibConstantsRcd.h" +#include "CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h" +#include "CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPedestalsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h" +#include "CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h" +#include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "HeterogeneousCore/CUDACore/interface/ConvertingESProducerT.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h" using EcalPedestalsGPUESProducer = ConvertingESProducerT; diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalMultifitParametersGPUESProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalMultifitParametersGPUESProducer.cc index 1743df5aa945d..406f5507be5fc 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalMultifitParametersGPUESProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalMultifitParametersGPUESProducer.cc @@ -2,6 +2,7 @@ #include #include +#include "CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h" #include "FWCore/Framework/interface/ESProducer.h" #include "FWCore/Framework/interface/ESProductHost.h" #include "FWCore/Framework/interface/ESTransientHandle.h" @@ -14,7 +15,6 @@ #include "FWCore/Utilities/interface/ReusableObjectHolder.h" #include "FWCore/Utilities/interface/typelookup.h" #include "HeterogeneousCore/CUDACore/interface/JobConfigurationGPURecord.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h" class EcalMultifitParametersGPUESProducer : public edm::ESProducer, public edm::EventSetupRecordIntervalFinder { public: diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.cu index 6e1b2a66c2507..8e77a58b66f4a 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.cu @@ -285,11 +285,9 @@ namespace ecal { // // Check for channels to be excluded from reconstruction // - // - // Default energy? Not to be updated if "ChannelStatusToBeExcluded" + // Default energy not to be updated if "ChannelStatusToBeExcluded" // Exploited later by the module "EcalRecHitConvertGPU2CPUFormat" - // - energy[inputCh] = -1; //---- AM: default, un-physical, ok + energy[inputCh] = -1; //un-physical default // truncate the chi2 if (chi2_in[inputCh] > 64) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.h index cb9c7f435d7b3..8d468a0f7f1ec 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitBuilderKernels.h @@ -9,7 +9,6 @@ #include "CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h" #include "DataFormats/Provenance/interface/Timestamp.h" -#include "Common.h" #include "DeclsForKernels.h" namespace ecal { diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitConvertGPU2CPUFormat.cc index 6df36f4a8b592..b71747a57db78 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitConvertGPU2CPUFormat.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitConvertGPU2CPUFormat.cc @@ -8,8 +8,6 @@ #include "FWCore/Framework/interface/stream/EDProducer.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" -#include "Common.h" - class EcalRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { public: explicit EcalRecHitConvertGPU2CPUFormat(edm::ParameterSet const& ps); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitParametersGPUESProducer.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitParametersGPUESProducer.cc index a63ed42cb2b70..8c6c8ce3ae236 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitParametersGPUESProducer.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitParametersGPUESProducer.cc @@ -2,6 +2,7 @@ #include #include +#include "CondFormats/EcalObjects/interface/EcalRecHitParametersGPU.h" #include "FWCore/Framework/interface/ESProducer.h" #include "FWCore/Framework/interface/ESProductHost.h" #include "FWCore/Framework/interface/ESTransientHandle.h" @@ -14,7 +15,6 @@ #include "FWCore/Utilities/interface/ReusableObjectHolder.h" #include "FWCore/Utilities/interface/typelookup.h" #include "HeterogeneousCore/CUDACore/interface/JobConfigurationGPURecord.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRecHitParametersGPU.h" class EcalRecHitParametersGPUESProducer : public edm::ESProducer, public edm::EventSetupRecordIntervalFinder { public: diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc index a6dabd37f8439..38c142f4c41f2 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalRecHitProducerGPU.cc @@ -9,6 +9,14 @@ #include "CondFormats/DataRecord/interface/EcalLaserAPDPNRatiosRefRcd.h" #include "CondFormats/DataRecord/interface/EcalLaserAlphasRcd.h" #include "CondFormats/DataRecord/interface/EcalLinearCorrectionsRcd.h" +#include "CondFormats/EcalObjects/interface/EcalIntercalibConstantsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAPDPNRatiosRefGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLaserAlphasGPU.h" +#include "CondFormats/EcalObjects/interface/EcalLinearCorrectionsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRecHitParametersGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitADCToGeVConstantGPU.h" +#include "CondFormats/EcalObjects/interface/EcalRechitChannelStatusGPU.h" #include "DataFormats/EcalRecHit/interface/EcalRecHit.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/EventSetup.h" @@ -20,14 +28,6 @@ #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalIntercalibConstantsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAPDPNRatiosRefGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLaserAlphasGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalLinearCorrectionsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRecHitParametersGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitADCToGeVConstantGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalRechitChannelStatusGPU.h" #include "EcalRecHitBuilderKernels.h" @@ -198,11 +198,11 @@ void EcalRecHitProducerGPU::acquire(edm::Event const& event, auto const& recHitParametersProduct = recHitParametersHandle_->getProduct(ctx.stream()); // set config ptrs : this is done to avoid changing things downstream - configParameters_.ChannelStatusToBeExcluded = recHitParametersProduct.ChannelStatusToBeExcluded; + configParameters_.ChannelStatusToBeExcluded = recHitParametersProduct.channelStatusToBeExcluded.get(); configParameters_.ChannelStatusToBeExcludedSize = std::get<0>(recHitParametersHandle_->getValues()).get().size(); - configParameters_.expanded_v_DB_reco_flags = recHitParametersProduct.expanded_v_DB_reco_flags; - configParameters_.expanded_Sizes_v_DB_reco_flags = recHitParametersProduct.expanded_Sizes_v_DB_reco_flags; - configParameters_.expanded_flagbit_v_DB_reco_flags = recHitParametersProduct.expanded_flagbit_v_DB_reco_flags; + configParameters_.expanded_v_DB_reco_flags = recHitParametersProduct.expanded_v_DB_reco_flags.get(); + configParameters_.expanded_Sizes_v_DB_reco_flags = recHitParametersProduct.expanded_Sizes_v_DB_reco_flags.get(); + configParameters_.expanded_flagbit_v_DB_reco_flags = recHitParametersProduct.expanded_flagbit_v_DB_reco_flags.get(); configParameters_.expanded_v_DB_reco_flagsSize = std::get<3>(recHitParametersHandle_->getValues()).get().size(); // bundle up conditions diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc index f7e57a61fdd96..b26fbe3a0c572 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitConvertGPU2CPUFormat.cc @@ -8,8 +8,6 @@ #include "FWCore/Framework/interface/stream/EDProducer.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" -#include "Common.h" - class EcalUncalibRecHitConvertGPU2CPUFormat : public edm::stream::EDProducer<> { public: explicit EcalUncalibRecHitConvertGPU2CPUFormat(edm::ParameterSet const& ps); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitMultiFitAlgoGPU.cu b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitMultiFitAlgoGPU.cu index 9d5a8a2ad1bd3..6d4f0b9a24220 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitMultiFitAlgoGPU.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitMultiFitAlgoGPU.cu @@ -15,7 +15,6 @@ #include "AmplitudeComputationCommonKernels.h" #include "AmplitudeComputationKernels.h" -#include "Common.h" #include "EcalUncalibRecHitMultiFitAlgoGPU.h" #include "TimeComputationKernels.h" @@ -262,10 +261,10 @@ namespace ecal { eventInputGPU.ebDigis.ids.get(), eventInputGPU.eeDigis.data.get(), eventInputGPU.eeDigis.ids.get(), - conditions.timeBiasCorrections.EBTimeCorrAmplitudeBins, - conditions.timeBiasCorrections.EETimeCorrAmplitudeBins, - conditions.timeBiasCorrections.EBTimeCorrShiftBins, - conditions.timeBiasCorrections.EETimeCorrShiftBins, + conditions.timeBiasCorrections.ebTimeCorrAmplitudeBins, + conditions.timeBiasCorrections.eeTimeCorrAmplitudeBins, + conditions.timeBiasCorrections.ebTimeCorrShiftBins, + conditions.timeBiasCorrections.eeTimeCorrShiftBins, scratch.timeMax.get(), scratch.timeError.get(), conditions.pedestals.rms_x12, @@ -276,8 +275,8 @@ namespace ecal { eventOutputGPU.recHitsEE.jitterError.get(), eventOutputGPU.recHitsEB.flags.get(), eventOutputGPU.recHitsEE.flags.get(), - conditions.timeBiasCorrections.EBTimeCorrAmplitudeBinsSize, - conditions.timeBiasCorrections.EETimeCorrAmplitudeBinsSize, + conditions.timeBiasCorrections.ebTimeCorrAmplitudeBinsSize, + conditions.timeBiasCorrections.eeTimeCorrAmplitudeBinsSize, configParameters.timeConstantTermEB, configParameters.timeConstantTermEE, conditions.timeOffsetConstant.getEBValue(), diff --git a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc index a321f35144c39..3729bddde895f 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc +++ b/RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitProducerGPU.cc @@ -8,6 +8,14 @@ #include "CondFormats/DataRecord/interface/EcalTimeBiasCorrectionsRcd.h" #include "CondFormats/DataRecord/interface/EcalTimeCalibConstantsRcd.h" #include "CondFormats/DataRecord/interface/EcalTimeOffsetConstantRcd.h" +#include "CondFormats/EcalObjects/interface/EcalGainRatiosGPU.h" +#include "CondFormats/EcalObjects/interface/EcalMultifitParametersGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPedestalsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseCovariancesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalPulseShapesGPU.h" +#include "CondFormats/EcalObjects/interface/EcalSamplesCorrelationGPU.h" +#include "CondFormats/EcalObjects/interface/EcalTimeBiasCorrectionsGPU.h" +#include "CondFormats/EcalObjects/interface/EcalTimeCalibConstantsGPU.h" #include "CondFormats/EcalObjects/interface/EcalTimeOffsetConstant.h" #include "DataFormats/EcalDigi/interface/EcalDigiCollections.h" #include "FWCore/Framework/interface/Event.h" @@ -18,16 +26,7 @@ #include "HeterogeneousCore/CUDACore/interface/JobConfigurationGPURecord.h" #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalGainRatiosGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalMultifitParametersGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPedestalsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseCovariancesGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalPulseShapesGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalSamplesCorrelationGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeBiasCorrectionsGPU.h" -#include "RecoLocalCalo/EcalRecAlgos/interface/EcalTimeCalibConstantsGPU.h" - -#include "Common.h" + #include "DeclsForKernels.h" #include "EcalUncalibRecHitMultiFitAlgoGPU.h" @@ -229,10 +228,10 @@ void EcalUncalibRecHitProducerGPU::acquire(edm::Event const& event, auto const& multifitParameters = multifitParametersData.getProduct(ctx.stream()); // assign ptrs/values: this is done not to change how things look downstream - configParameters_.amplitudeFitParametersEB = multifitParameters.amplitudeFitParametersEB; - configParameters_.amplitudeFitParametersEE = multifitParameters.amplitudeFitParametersEE; - configParameters_.timeFitParametersEB = multifitParameters.timeFitParametersEB; - configParameters_.timeFitParametersEE = multifitParameters.timeFitParametersEE; + configParameters_.amplitudeFitParametersEB = multifitParameters.amplitudeFitParametersEB.get(); + configParameters_.amplitudeFitParametersEE = multifitParameters.amplitudeFitParametersEE.get(); + configParameters_.timeFitParametersEB = multifitParameters.timeFitParametersEB.get(); + configParameters_.timeFitParametersEE = multifitParameters.timeFitParametersEE.get(); configParameters_.timeFitParametersSizeEB = multifitParametersData.getValues()[2].get().size(); configParameters_.timeFitParametersSizeEE = multifitParametersData.getValues()[3].get().size(); diff --git a/RecoLocalCalo/EcalRecProducers/plugins/KernelHelpers.cu b/RecoLocalCalo/EcalRecProducers/plugins/KernelHelpers.cu index 5316ed87d6ecc..a05e69f7b0442 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/KernelHelpers.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/KernelHelpers.cu @@ -47,10 +47,6 @@ namespace ecal { return dccFromSm(ism); } - // - // ---- why on hell things are so complex and not simple ??? - // - __device__ int lm_channel(int iX, int iY) { static const int idx_[] = { // clang-format off diff --git a/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.cu b/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.cu index 9c2d2fc986c08..e67802f44c8df 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.cu +++ b/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.cu @@ -4,12 +4,12 @@ #include #include "DataFormats/EcalDigi/interface/EcalDataFrame.h" +#include "DataFormats/EcalDigi/interface/EcalMGPASample.h" #include "DataFormats/EcalRecHit/interface/EcalUncalibratedRecHit.h" #include "DataFormats/Math/interface/approx_exp.h" #include "DataFormats/Math/interface/approx_log.h" #include "FWCore/Utilities/interface/CMSUnrollLoop.h" -#include "Common.h" #include "TimeComputationKernels.h" #include "KernelHelpers.h" @@ -693,8 +693,8 @@ namespace ecal { if (!use_sample(sample_mask, sample)) return; - const auto gainIdPrev = ecal::mgpa::gainId(digis[inputGtx - 1]); - const auto gainIdNext = ecal::mgpa::gainId(digis[inputGtx]); + const auto gainIdPrev = ecalMGPA::gainId(digis[inputGtx - 1]); + const auto gainIdNext = ecalMGPA::gainId(digis[inputGtx]); if (gainIdPrev >= 1 && gainIdPrev <= 3 && gainIdNext >= 1 && gainIdNext <= 3 && gainIdPrev < gainIdNext) { sample_values[gtx - 1] = 0; sample_value_errors[gtx - 1] = 1e+9; @@ -849,10 +849,10 @@ namespace ecal { ScalarType* shrSampleValueErrors = shrSampleValues + blockDim.x; // 0 and 1 sample values - const auto adc0 = ecal::mgpa::adc(digis[input_ch_start]); - const auto gainId0 = ecal::mgpa::gainId(digis[input_ch_start]); - const auto adc1 = ecal::mgpa::adc(digis[input_ch_start + 1]); - const auto gainId1 = ecal::mgpa::gainId(digis[input_ch_start + 1]); + const auto adc0 = ecalMGPA::adc(digis[input_ch_start]); + const auto gainId0 = ecalMGPA::gainId(digis[input_ch_start]); + const auto adc1 = ecalMGPA::adc(digis[input_ch_start + 1]); + const auto gainId1 = ecalMGPA::gainId(digis[input_ch_start + 1]); const auto did = DetId{dids[inputCh]}; const auto isBarrel = did.subdetId() == EcalBarrel; const auto sample_mask = did.subdetId() == EcalBarrel ? sample_maskEB : sample_maskEE; @@ -875,8 +875,8 @@ namespace ecal { } // ped subtracted and gain-renormalized samples. - const auto gainId = ecal::mgpa::gainId(digis[inputTx]); - const auto adc = ecal::mgpa::adc(digis[inputTx]); + const auto gainId = ecalMGPA::gainId(digis[inputTx]); + const auto adc = ecalMGPA::adc(digis[inputTx]); bool bad = false; SampleVector::Scalar sample_value, sample_value_error; @@ -1112,7 +1112,7 @@ namespace ecal { auto threshM = outOfTimeThreshG12m; if (amplitude > 3000.) { for (int isample = 0; isample < nsamples; isample++) { - int gainid = ecal::mgpa::gainId(digis[nsamples * inputGtx + isample]); + int gainid = ecalMGPA::gainId(digis[nsamples * inputGtx + isample]); if (gainid != 1) { threshP = outOfTimeThreshG61p; threshM = outOfTimeThreshG61m; diff --git a/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.h index a9b1c69678abd..dea6bad26fa0d 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.h +++ b/RecoLocalCalo/EcalRecProducers/plugins/TimeComputationKernels.h @@ -9,7 +9,6 @@ #include "DataFormats/Math/interface/approx_exp.h" #include "DataFormats/Math/interface/approx_log.h" -#include "Common.h" #include "DeclsForKernels.h" #include "EigenMatrixTypes_gpu.h"