diff --git a/UserCode/CodeGPU/plugins/HGCalRecHitKernelImpl.cc b/UserCode/CodeGPU/plugins/HGCalRecHitKernelImpl.cc deleted file mode 100644 index 93d4f49415fca..0000000000000 --- a/UserCode/CodeGPU/plugins/HGCalRecHitKernelImpl.cc +++ /dev/null @@ -1,11 +0,0 @@ - - -__global__ -void HGCeeRecHitKernel(HGCUncalibratedRecHit* oldhits, - HGCUncalibratedRecHit* newhits, size_t length) -{ - for (size_t i = blockDim.x * blockIdx.x + threadIdx.x; i < length; i += blockDim.x * gridDim.x) - { - newhits[i] = oldhits[i]; - } -} diff --git a/UserCode/CodeGPU/plugins/HGCalRecHitKernelImpl.h b/UserCode/CodeGPU/plugins/HGCalRecHitKernelImpl.h deleted file mode 100644 index 029d3530cc959..0000000000000 --- a/UserCode/CodeGPU/plugins/HGCalRecHitKernelImpl.h +++ /dev/null @@ -1,59 +0,0 @@ -#ifndef HGCalRecHitKernelImpl_cuh -#define HGCalRecHitKernelImpl_cuh - -#include "FWCore/Utilities/interface/Exception.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "Utils.h" - -#include //std::swap -#include - -class KernelManagerBase { -public: - explicit KernelManagerBase(){}; - virtual ~KernelManagerBase(){}; - - virtual void run_kernels() = 0; - -protected: - virtual void assign_and_transfer_to_device() = 0; - virtual void transfer_to_host_and_synchronize() = 0; - virtual void reuse_device_pointers() = 0; - - dim3 dimGrid; - dim3 dimBlock; -}; - -//the class assumes that the sizes of the arrays pointed to and the size of the collection are constant -class KernelManagerHGCalRecHit: private KernelManagerBase { -public: - explicit KernelManagerHGCalRecHit(const edm::SortedCollection&, const detectortype&); - ~KernelManagerHGCalRecHit(); - void run_kernels(); - const edm::SortedCollection& get_new_collection(); - -private: - void ee_step1_wrapper(); - void hef_step1_wrapper(); - void heb_step1_wrapper(); - void assign_and_transfer_to_device() override; - void transfer_to_host_and_synchronize() override; - void reuse_device_pointers() override; - - size_t shits_, sbytes_; - const detectortype dtype_; - edm::SortedCollection oldhits_collection_; - HGCUncalibratedRecHit *h_oldhits_, *h_newhits_; //host pointers - HGCUncalibratedRecHit *d_oldhits_, *d_newhits_; //device pointers -}; - -__global__ -void ee_step1(HGCUncalibratedRecHit *, HGCUncalibratedRecHit *, const size_t& length); -__global__ -void hef_step1(HGCUncalibratedRecHit *, HGCUncalibratedRecHit *, const size_t& length); -__global__ -void heb_step1(HGCUncalibratedRecHit *, HGCUncalibratedRecHit *, const size_t& length); - - -#endif //HGCalRecHitKernelImpl_cuh diff --git a/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.cc b/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.cc index e99dad2020465..0bb7b14d7fdd2 100644 --- a/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.cc +++ b/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.cc @@ -3,11 +3,6 @@ HeterogeneousHGCalEERecHitProducer::HeterogeneousHGCalEERecHitProducer(const edm::ParameterSet& ps): token_(consumes(ps.getParameter("HGCEEUncalibRecHitsTok"))) { - histo1_ = fs->make( "energy" , "E", 100, 0., 10. ); - histo2_ = fs->make( "time" , "t", 100, 0., 10. ); - histo3_ = fs->make( "timeError" , "time_error", 100, 0., 10. ); - histo4_ = fs->make( "son" , "son", 100, 0., 10. ); - nhitsmax_ = ps.getParameter("nhitsmax"); cdata_.hgcEE_keV2DIGI_ = ps.getParameter("HGCEE_keV2DIGI"); cdata_.xmin_ = ps.getParameter("minValSiPar"); //float @@ -65,7 +60,6 @@ void HeterogeneousHGCalEERecHitProducer::acquire(edm::Event const& event, edm::E const auto &hits_ee = *handle_ee_; unsigned int nhits = hits_ee.size(); - std::cout << "EE hits: " << nhits << std::endl; convert_collection_data_to_soa_(hits_ee, old_soa_, nhits); kmdata_ = new KernelModifiableData(nhitsmax_, stride_, old_soa_, d_oldhits_, d_newhits_, d_newhits_final_, h_newhits_); @@ -73,9 +67,7 @@ void HeterogeneousHGCalEERecHitProducer::acquire(edm::Event const& event, edm::E kernel_manager.run_kernels(h_kcdata_, d_kcdata_); new_soa_ = kernel_manager.get_output(); - //print_to_histograms(kmdata_->h_out, histo1_, histo2_, histo3_, histo4_, nhits); - - rechits_ = std::make_unique< HGCRecHitCollection >(); + rechits_ = std::make_unique(); convert_soa_data_to_collection_(*rechits_, new_soa_, nhits); } diff --git a/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.h b/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.h index 2aa5350908cde..23683bdce4392 100644 --- a/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.h +++ b/UserCode/CodeGPU/plugins/HeterogeneousHGCalEERecHitProducer.h @@ -34,7 +34,7 @@ #include "HeterogeneousHGCalProducerMemoryWrapper.h" #include "KernelManager.h" - +#include "Utils.h" class HeterogeneousHGCalEERecHitProducer: public edm::stream::EDProducer { diff --git a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.cc b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.cc index 835140da3143e..660ef13e002cd 100644 --- a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.cc +++ b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.cc @@ -3,11 +3,6 @@ HeterogeneousHGCalHEBRecHitProducer::HeterogeneousHGCalHEBRecHitProducer(const edm::ParameterSet& ps): token_(consumes(ps.getParameter("HGCHEBUncalibRecHitsTok"))) { - histo1_ = fs->make( "energy" , "E", 100, 0., 10. ); - histo2_ = fs->make( "time" , "t", 100, 0., 10. ); - histo3_ = fs->make( "timeError" , "time_error", 100, 0., 10. ); - histo4_ = fs->make( "son" , "son", 100, 0., 10. ); - nhitsmax_ = ps.getParameter("nhitsmax"); cdata_.hgcHEB_keV2DIGI_ = ps.getParameter("HGCHEB_keV2DIGI"); cdata_.hgcHEB_noise_MIP_ = ps.getParameter("HGCHEB_noise_MIP").getParameter("noise_MIP"); @@ -48,13 +43,11 @@ HeterogeneousHGCalHEBRecHitProducer::~HeterogeneousHGCalHEBRecHitProducer() void HeterogeneousHGCalHEBRecHitProducer::acquire(edm::Event const& event, edm::EventSetup const& setup, edm::WaitingTaskWithArenaHolder w) { const cms::cuda::ScopedContextAcquire ctx{event.streamID(), std::move(w), ctxState_}; - set_geometry_(setup); event.getByToken(token_, handle_heb_); const auto &hits_heb = *handle_heb_; unsigned int nhits = hits_heb.size(); - std::cout << "HEB hits: " << nhits << std::endl; convert_collection_data_to_soa_(hits_heb, old_soa_, nhits); kmdata_ = new KernelModifiableData(nhitsmax_, stride_, old_soa_, d_oldhits_, d_newhits_, d_newhits_final_, h_newhits_); @@ -62,9 +55,7 @@ void HeterogeneousHGCalHEBRecHitProducer::acquire(edm::Event const& event, edm:: kernel_manager.run_kernels(h_kcdata_, d_kcdata_); new_soa_ = kernel_manager.get_output(); - //print_to_histograms(kmdata_->h_out, histo1_, histo2_, histo3_, histo4_, nhits); - - rechits_ = std::make_unique< HGCRecHitCollection >(); + rechits_ = std::make_unique(); convert_soa_data_to_collection_(*rechits_, new_soa_, nhits); } diff --git a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.h b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.h index 13ad5736395c5..3a591887447f9 100644 --- a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.h +++ b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEBRecHitProducer.h @@ -34,7 +34,7 @@ #include "HeterogeneousHGCalProducerMemoryWrapper.h" #include "KernelManager.h" - +#include "Utils.h" class HeterogeneousHGCalHEBRecHitProducer: public edm::stream::EDProducer { diff --git a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.cc b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.cc index e8c9f85290b19..5dc349da4b24f 100644 --- a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.cc +++ b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.cc @@ -3,12 +3,6 @@ HeterogeneousHGCalHEFRecHitProducer::HeterogeneousHGCalHEFRecHitProducer(const edm::ParameterSet& ps): token_(consumes(ps.getParameter("HGCHEFUncalibRecHitsTok"))) { - //histogram drawing for debugging: to be removed at a later stage - histo1_ = fs->make( "energy" , "E", 100, 0., .2 ); - histo2_ = fs->make( "time" , "t", 100, 0., 10. ); - histo3_ = fs->make( "timeError" , "time_error", 100, 0., 10. ); - histo4_ = fs->make( "son" , "son", 32, 0., 31. ); - nhitsmax_ = ps.getParameter("nhitsmax"); cdata_.hgcHEF_keV2DIGI_ = ps.getParameter("HGCHEF_keV2DIGI"); cdata_.xmin_ = ps.getParameter("minValSiPar"); //float @@ -67,17 +61,14 @@ void HeterogeneousHGCalHEFRecHitProducer::acquire(edm::Event const& event, edm:: const auto &hits_hef = *handle_hef_; unsigned int nhits = hits_hef.size(); - std::cout << "HEF hits: " << nhits << std::endl; convert_collection_data_to_soa_(hits_hef, old_soa_, nhits); kmdata_ = new KernelModifiableData(nhitsmax_, stride_, old_soa_, d_oldhits_, d_newhits_, d_newhits_final_, h_newhits_); KernelManagerHGCalRecHit kernel_manager(kmdata_); kernel_manager.run_kernels(h_kcdata_, d_kcdata_); new_soa_ = kernel_manager.get_output(); - - //print_to_histograms(kmdata_->h_out, histo1_, histo2_, histo3_, histo4_, nhits); - - rechits_ = std::make_unique< HGCRecHitCollection >(); + + rechits_ = std::make_unique(); convert_soa_data_to_collection_(*rechits_, new_soa_, nhits); } diff --git a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.h b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.h index 94855ac50b325..109cb04089eaa 100644 --- a/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.h +++ b/UserCode/CodeGPU/plugins/HeterogeneousHGCalHEFRecHitProducer.h @@ -34,7 +34,7 @@ #include "HeterogeneousHGCalProducerMemoryWrapper.h" #include "KernelManager.h" - +#include "Utils.h" class HeterogeneousHGCalHEFRecHitProducer: public edm::stream::EDProducer {