SiPixelGainForHLTonGPU.h

Go to the documentation of this file.

#ifndef CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h
#define CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h

#include <cstdint>
#include <cstdio>
#include <tuple>

// including <cuda_runtime.h> would pull in the dependency on all of CUDA;
// instead, just define away the CUDA specific attributes to keep GCC happy.
#ifndef __CUDACC__
#ifndef __host__
#define __host__
#endif  // __host__
#ifndef __device__
#define __device__
#endif  // __device__
#endif  // __CUDACC__

#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"

struct SiPixelGainForHLTonGPU_DecodingStructure {
  uint8_t gain;
  uint8_t ped;
};

// copy of SiPixelGainCalibrationForHLT
class SiPixelGainForHLTonGPU {
public:
  using DecodingStructure = SiPixelGainForHLTonGPU_DecodingStructure;

  using Range = std::pair<uint32_t, uint32_t>;

  inline __host__ __device__ std::pair<float, float> getPedAndGain(
      uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn) const {
    auto range = rangeAndCols_[moduleInd].first;
    auto nCols = rangeAndCols_[moduleInd].second;

    // 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
    unsigned int lengthOfColumnData = (range.second - range.first) / nCols;
    unsigned int lengthOfAveragedDataInEachColumn = 2;  // we always only have two values per column averaged block
    unsigned int numberOfDataBlocksToSkip = row / numberOfRowsAveragedOver_;

    auto offset = range.first + col * lengthOfColumnData + lengthOfAveragedDataInEachColumn * numberOfDataBlocksToSkip;

    assert(offset < range.second);
    assert(offset < 3088384);
    assert(0 == offset % 2);

    DecodingStructure const* __restrict__ lp = v_pedestals_;
    auto s = lp[offset / 2];

    isDeadColumn = (s.ped & 0xFF) == deadFlag_;
    isNoisyColumn = (s.ped & 0xFF) == noisyFlag_;

    return std::make_pair(decodePed(s.ped & 0xFF), decodeGain(s.gain & 0xFF));
  }

  constexpr float decodeGain(unsigned int gain) const { return float(gain) * gainPrecision_ + minGain_; }
  constexpr float decodePed(unsigned int ped) const { return float(ped) * pedPrecision_ + minPed_; }

  DecodingStructure* v_pedestals_;
  std::pair<Range, int> rangeAndCols_[gpuClustering::maxNumModules];

  float minPed_, maxPed_, minGain_, maxGain_;
  float pedPrecision_, gainPrecision_;

  unsigned int numberOfRowsAveragedOver_;  // this is 80!!!!
  unsigned int nBinsToUseForEncoding_;
  unsigned int deadFlag_;
  unsigned int noisyFlag_;
};

#endif  // CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h