Warning, /RecoLocalCalo/EcalRecProducers/plugins/EcalUncalibRecHitPhase2WeightsKernels.cu is written in an unsupported language. File is not indexed.
0001 #include <cuda.h>
0002
0003 #include "FWCore/Utilities/interface/CMSUnrollLoop.h"
0004 #include "DataFormats/EcalRecHit/interface/EcalUncalibratedRecHit.h"
0005 #include "DataFormats/EcalDigi/interface/EcalLiteDTUSample.h"
0006 #include "DataFormats/EcalDigi/interface/EcalConstants.h"
0007
0008 #include "EcalUncalibRecHitPhase2WeightsKernels.h"
0009
0010 namespace ecal {
0011 namespace weights {
0012
0013 __global__ void Phase2WeightsKernel(uint16_t const* digis_in,
0014 uint32_t const* __restrict__ dids,
0015 ::ecal::reco::StorageScalarType* __restrict__ amplitude,
0016 ::ecal::reco::StorageScalarType* __restrict__ amplitudeError,
0017 uint32_t* __restrict__ dids_out,
0018 int const nchannels,
0019 double const* __restrict__ weights,
0020 uint32_t* __restrict__ flags) {
0021 constexpr int nsamples = EcalDataFrame_Ph2::MAXSAMPLES;
0022 unsigned int nchannels_per_block = blockDim.x;
0023
0024 // copy data from global to shared memory
0025 extern __shared__ char shared_mem[];
0026 double* shr_weights = reinterpret_cast<double*>(shared_mem); // nsamples elements
0027 float* shr_amp = reinterpret_cast<float*>(shr_weights + nsamples); // nchannels_per_block elements
0028 uint16_t* shr_digis = reinterpret_cast<uint16_t*>(shr_amp + nchannels_per_block); // nchannels_per_block elements
0029 for (int i = 0; i < nsamples; ++i)
0030 shr_weights[i] = weights[i];
0031
0032 unsigned int const threadx = threadIdx.x;
0033 unsigned int const blockx = blockIdx.x;
0034
0035 for (int sample = 0; sample < nsamples; ++sample) {
0036 int const idx = threadx * nsamples + sample;
0037 shr_digis[idx] = digis_in[blockx * nchannels_per_block * nsamples + idx];
0038 }
0039 shr_amp[threadx] = 0.;
0040
0041 __syncthreads();
0042
0043 const auto first = threadIdx.x + blockIdx.x * blockDim.x;
0044 const auto stride = blockDim.x * gridDim.x;
0045 for (auto tx = first; tx < nchannels; tx += stride) {
0046 auto const did = DetId{dids[tx]};
0047 CMS_UNROLL_LOOP
0048 for (int sample = 0; sample < nsamples; ++sample) {
0049 const unsigned int idx = threadx * nsamples + sample;
0050 const auto shr_digi = shr_digis[idx];
0051 shr_amp[threadx] += (static_cast<float>(ecalLiteDTU::adc(shr_digi)) *
0052 ecalPh2::gains[ecalLiteDTU::gainId(shr_digi)] * shr_weights[sample]);
0053 }
0054 const unsigned int tdx = threadx * nsamples;
0055 amplitude[tx] = shr_amp[threadx];
0056 amplitudeError[tx] = 1.0f;
0057 dids_out[tx] = did.rawId();
0058 flags[tx] = 0;
0059 if (ecalLiteDTU::gainId(shr_digis[tdx + nsamples - 1])) {
0060 flags[tx] = 0x1 << EcalUncalibratedRecHit::kHasSwitchToGain1;
0061 }
0062 } //if within nchannels
0063 } //kernel
0064 } //namespace weights
0065 } //namespace ecal
