Warning, /RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsAlgoGPU.cu is written in an unsupported language. File is not indexed.
0001 #include "DataFormats/EcalDigi/interface/EcalDataFrame_Ph2.h"
0002 #include "DataFormats/EcalDigi/interface/EcalDigiCollections.h"
0003
0004 #include "EcalUncalibRecHitPhase2WeightsKernels.h"
0005 #include "EcalUncalibRecHitPhase2WeightsAlgoGPU.h"
0006
0007 namespace ecal {
0008 namespace weights {
0009
0010 void phase2Weights(ecal::DigisCollection<calo::common::DevStoragePolicy> const& digis,
0011 EventOutputDataGPU& eventOutputGPU,
0012 cms::cuda::device::unique_ptr<double[]>& weights_d,
0013 cudaStream_t cudaStream) {
0014 unsigned int const totalChannels = digis.size;
0015 // 64 threads per block best occupancy from Nsight compute profiler
0016 unsigned int const threads_1d = 64;
0017 unsigned int const blocks_1d = (totalChannels + threads_1d - 1) / threads_1d;
0018 int shared_bytes = EcalDataFrame_Ph2::MAXSAMPLES * sizeof(double) +
0019 threads_1d * (EcalDataFrame_Ph2::MAXSAMPLES * (sizeof(uint16_t)) + sizeof(float));
0020 Phase2WeightsKernel<<<blocks_1d, threads_1d, shared_bytes, cudaStream>>>(
0021 digis.data.get(),
0022 digis.ids.get(),
0023 eventOutputGPU.recHits.amplitude.get(),
0024 eventOutputGPU.recHits.amplitudeError.get(),
0025 eventOutputGPU.recHits.did.get(),
0026 totalChannels,
0027 weights_d.get(),
0028 eventOutputGPU.recHits.flags.get());
0029 cudaCheck(cudaGetLastError());
0030 }
0031
0032 } // namespace weights
0033 } // namespace ecal
