HGCalRecProducers/plugins/KernelManagerHGCalRecHit.cu

0001 #include <cuda.h>
0002 #include <cuda_runtime.h>
0003 #include <inttypes.h>
0004 #include "RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.h"
0005 #include "RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cuh"
0006
0007 namespace {  //kernel parameters
0008   dim3 nb_rechits_;
0009   constexpr dim3 nt_rechits_(1024);
0010 }  // namespace
0011
0012 KernelManagerHGCalRecHit::KernelManagerHGCalRecHit(const HGCUncalibRecHitSoA& h_uncalibSoA,
0013                                                    const HGCUncalibRecHitSoA& d_uncalibSoA,
0014                                                    const HGCRecHitSoA& d_calibSoA)
0015     : h_uncalibSoA_(h_uncalibSoA), d_uncalibSoA_(d_uncalibSoA), d_calibSoA_(d_calibSoA) {
0016   nhits_ = h_uncalibSoA_.nhits_;
0017   pad_ = h_uncalibSoA_.pad_;
0018   ::nb_rechits_ = (pad_ + ::nt_rechits_.x - 1) / ::nt_rechits_.x;
0019   nbytes_device_ = d_uncalibSoA_.nbytes_ * pad_;
0020 }
0021
0022 KernelManagerHGCalRecHit::KernelManagerHGCalRecHit(const HGCRecHitSoA& h_calibSoA_,
0023                                                    const ConstHGCRecHitSoA& d_calibConstSoA)
0024     : h_calibSoA_(h_calibSoA_), d_calibConstSoA_(d_calibConstSoA) {
0025   nhits_ = h_calibSoA_.nhits_;
0026   pad_ = h_calibSoA_.pad_;
0027   ::nb_rechits_ = (pad_ + ::nt_rechits_.x - 1) / ::nt_rechits_.x;
0028   nbytes_host_ = h_calibSoA_.nbytes_ * pad_;
0029 }
0030
0031 KernelManagerHGCalRecHit::~KernelManagerHGCalRecHit() {}
0032
0033 void KernelManagerHGCalRecHit::transfer_soa_to_device_(const cudaStream_t& stream) {
0034   cudaCheck(cudaMemcpyAsync(
0035       d_uncalibSoA_.amplitude_, h_uncalibSoA_.amplitude_, nbytes_device_, cudaMemcpyHostToDevice, stream));
0036 }
0037
0038 void KernelManagerHGCalRecHit::transfer_soa_to_host(const cudaStream_t& stream) {
0039   cudaCheck(
0040       cudaMemcpyAsync(h_calibSoA_.energy_, d_calibConstSoA_.energy_, nbytes_host_, cudaMemcpyDeviceToHost, stream));
0041 }
0042
0043 void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGCeeUncalibRecHitConstantData>* kcdata,
0044                                            const cudaStream_t& stream) {
0045   transfer_soa_to_device_(stream);
0046   ee_to_rechit<<<::nb_rechits_, ::nt_rechits_, 0, stream>>>(d_calibSoA_, d_uncalibSoA_, kcdata->data_, nhits_);
0047   cudaCheck(cudaGetLastError());
0048 }
0049
0050 void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChefUncalibRecHitConstantData>* kcdata,
0051                                            const cudaStream_t& stream) {
0052   transfer_soa_to_device_(stream);
0053   hef_to_rechit<<<::nb_rechits_, ::nt_rechits_, 0, stream>>>(d_calibSoA_, d_uncalibSoA_, kcdata->data_, nhits_);
0054   cudaCheck(cudaGetLastError());
0055 }
0056
0057 void KernelManagerHGCalRecHit::run_kernels(const KernelConstantData<HGChebUncalibRecHitConstantData>* kcdata,
0058                                            const cudaStream_t& stream) {
0059   transfer_soa_to_device_(stream);
0060   heb_to_rechit<<<::nb_rechits_, ::nt_rechits_, 0, stream>>>(d_calibSoA_, d_uncalibSoA_, kcdata->data_, nhits_);
0061   cudaCheck(cudaGetLastError());
0062 }