SiPixelObjects/interface/SiPixelGainForHLTonGPU.h

0001 #ifndef CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h
0002 #define CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h
0003
0004 #include <cstdint>
0005 #include <cstdio>
0006 #include <tuple>
0007
0008 // including <cuda_runtime.h> would pull in the dependency on all of CUDA;
0009 // instead, just define away the CUDA specific attributes to keep GCC happy.
0010 #ifndef __CUDACC__
0011 #ifndef __host__
0012 #define __host__
0013 #endif  // __host__
0014 #ifndef __device__
0015 #define __device__
0016 #endif  // __device__
0017 #endif  // __CUDACC__
0018
0019 #include "Geometry/CommonTopologies/interface/SimplePixelTopology.h"
0020 #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
0021
0022 struct SiPixelGainForHLTonGPU_DecodingStructure {
0023   uint8_t gain;
0024   uint8_t ped;
0025 };
0026
0027 // copy of SiPixelGainCalibrationForHLT
0028 class SiPixelGainForHLTonGPU {
0029 public:
0030   using DecodingStructure = SiPixelGainForHLTonGPU_DecodingStructure;
0031
0032   using Range = std::pair<uint32_t, uint32_t>;
0033
0034   inline __host__ __device__ std::pair<float, float> getPedAndGain(
0035       uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn) const {
0036     auto range = rangeAndCols_[moduleInd].first;
0037     auto nCols = rangeAndCols_[moduleInd].second;
0038     // determine what averaged data block we are in (there should be 1 or 2 of these depending on if plaquette is 1 by X or 2 by X
0039     unsigned int lengthOfColumnData = (range.second - range.first) / nCols;
0040     unsigned int lengthOfAveragedDataInEachColumn = 2;  // we always only have two values per column averaged block
0041     unsigned int numberOfDataBlocksToSkip = row / numberOfRowsAveragedOver_;
0042
0043     auto offset = range.first + col * lengthOfColumnData + lengthOfAveragedDataInEachColumn * numberOfDataBlocksToSkip;
0044
0045     assert(offset < range.second);
0046     assert(offset < 3088384);
0047     assert(0 == offset % 2);
0048
0049     DecodingStructure const* __restrict__ lp = v_pedestals_;
0050     auto s = lp[offset / 2];
0051
0052     isDeadColumn = (s.ped & 0xFF) == deadFlag_;
0053     isNoisyColumn = (s.ped & 0xFF) == noisyFlag_;
0054
0055     return std::make_pair(decodePed(s.ped & 0xFF), decodeGain(s.gain & 0xFF));
0056   }
0057
0058   constexpr float decodeGain(unsigned int gain) const { return float(gain) * gainPrecision_ + minGain_; }
0059   constexpr float decodePed(unsigned int ped) const { return float(ped) * pedPrecision_ + minPed_; }
0060
0061   DecodingStructure* v_pedestals_;
0062   std::pair<Range, int> rangeAndCols_[phase1PixelTopology::numberOfModules];
0063
0064   float minPed_, maxPed_, minGain_, maxGain_;
0065   float pedPrecision_, gainPrecision_;
0066
0067   unsigned int numberOfRowsAveragedOver_;  // this is 80!!!!
0068   unsigned int nBinsToUseForEncoding_;
0069   unsigned int deadFlag_;
0070   unsigned int noisyFlag_;
0071 };
0072
0073 #endif  // CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h