gpuCalibPixel Namespace Reference

Functions
template<bool isRun2>
__global__ void	calibDigis (uint16_t *id, uint16_t const *__restrict__ x, uint16_t const *__restrict__ y, uint16_t *adc, SiPixelGainForHLTonGPU const *__restrict__ ped, int numElements, uint32_t *__restrict__ moduleStart, uint32_t *__restrict__ nClustersInModule, uint32_t *__restrict__ clusModuleStart)
__global__ void	calibDigisPhase2 (uint16_t *id, uint16_t *adc, int numElements, uint32_t *__restrict__ moduleStart, uint32_t *__restrict__ nClustersInModule, uint32_t *__restrict__ clusModuleStart)

Variables
constexpr float	ElectronPerADCGain = 600
constexpr uint16_t	Phase2DigiBaseline = 1500
constexpr uint8_t	Phase2KinkADC = 8
constexpr int8_t	Phase2ReadoutMode = 3
constexpr int	VCalChargeThreshold = 100
constexpr float	VCaltoElectronGain = 47
constexpr float	VCaltoElectronGain_L1 = 50
constexpr float	VCaltoElectronOffset = -60
constexpr float	VCaltoElectronOffset_L1 = -670

Function Documentation

calibDigis()

template<bool isRun2>

__global__ void gpuCalibPixel::calibDigis	(	uint16_t *	id,
		uint16_t const *__restrict__	x,
		uint16_t const *__restrict__	y,
		uint16_t *	adc,
		SiPixelGainForHLTonGPU const *__restrict__	ped,
		int	numElements,
		uint32_t *__restrict__	moduleStart,
		uint32_t *__restrict__	nClustersInModule,
		uint32_t *__restrict__	clusModuleStart
	)

Definition at line 32 of file gpuCalibPixel.h.

References gpuClustering::adc, cms::cudacompat::blockDim, cms::cudacompat::blockIdx, cuy::col, first, dqmMemoryStats::float, PedestalClient_cfi::gain, cms::cudacompat::gridDim, mps_fire::i, gpuClustering::invalidModuleId, geometryPPS_CMSxz_fromDD_2016_cfi::isRun2, SiStripPI::max, gpuClustering::moduleStart, gpuClustering::nClustersInModule, phase1PixelTopology::numberOfModules, gpuClustering::numElements, hltrates_dqm_sourceclient-live_cfg::offset, EcalCondDBWriter_cfi::pedestal, runTheMatrix::ret, cms::cudacompat::threadIdx, VCaltoElectronGain, VCaltoElectronGain_L1, VCaltoElectronOffset, VCaltoElectronOffset_L1, and x.

    {
    int first = blockDim.x * blockIdx.x + threadIdx.x;

    // zero for next kernels...
    if (0 == first)
      clusModuleStart[0] = moduleStart[0] = 0;
    for (int i = first; i < phase1PixelTopology::numberOfModules; i += gridDim.x * blockDim.x) {
      nClustersInModule[i] = 0;
    }

    for (int i = first; i < numElements; i += gridDim.x * blockDim.x) {
      if (invalidModuleId == id[i])
        continue;

      bool isDeadColumn = false, isNoisyColumn = false;

      int row = x[i];
      int col = y[i];
      auto ret = ped->getPedAndGain(id[i], col, row, isDeadColumn, isNoisyColumn);
      float pedestal = ret.first;
      float gain = ret.second;
      // float pedestal = 0; float gain = 1.;
      if (isDeadColumn | isNoisyColumn) {
        printf("bad pixel at %d in %d\n", i, id[i]);
        id[i] = invalidModuleId;
        adc[i] = 0;
      } else {
        float vcal = float(adc[i]) * gain - pedestal * gain;
        if constexpr (isRun2) {
          float conversionFactor = id[i] < 96 ? VCaltoElectronGain_L1 : VCaltoElectronGain;
          float offset = id[i] < 96 ? VCaltoElectronOffset_L1 : VCaltoElectronOffset;
          vcal = vcal * conversionFactor + offset;
        }
        adc[i] = std::clamp(int(vcal), 100, int(std::numeric_limits<uint16_t>::max()));
      }
    }
  }

calibDigisPhase2()

__global__ void gpuCalibPixel::calibDigisPhase2	(	uint16_t *	id,
		uint16_t *	adc,
		int	numElements,
		uint32_t *__restrict__	moduleStart,
		uint32_t *__restrict__	nClustersInModule,
		uint32_t *__restrict__	clusModuleStart
	)

Definition at line 79 of file gpuCalibPixel.h.

References gpuClustering::adc, cms::cudacompat::blockDim, cms::cudacompat::blockIdx, ElectronPerADCGain, first, cms::cudacompat::gridDim, mps_fire::i, createfilelist::int, gpuClustering::invalidModuleId, SiStripPI::max, SiStripPI::min, ALCARECOPromptCalibProdSiPixelAli0T_cff::mode, gpuClustering::moduleStart, gpuClustering::nClustersInModule, phase2PixelTopology::numberOfModules, gpuClustering::numElements, Phase2DigiBaseline, Phase2KinkADC, Phase2ReadoutMode, and cms::cudacompat::threadIdx.

    {
    int first = blockDim.x * blockIdx.x + threadIdx.x;
    // zero for next kernels...

    if (0 == first)
      clusModuleStart[0] = moduleStart[0] = 0;
    for (int i = first; i < phase2PixelTopology::numberOfModules; i += gridDim.x * blockDim.x) {
      nClustersInModule[i] = 0;
    }

    for (int i = first; i < numElements; i += gridDim.x * blockDim.x) {
      if (invalidModuleId == id[i])
        continue;

      constexpr int mode = (Phase2ReadoutMode < -1 ? -1 : Phase2ReadoutMode);

      int adc_int = adc[i];

      if constexpr (mode < 0)
        adc_int = int(adc_int * ElectronPerADCGain);
      else {
        if (adc_int < Phase2KinkADC)
          adc_int = int((adc_int - 0.5) * ElectronPerADCGain);
        else {
          constexpr int8_t dspp = (Phase2ReadoutMode < 10 ? Phase2ReadoutMode : 10);
          constexpr int8_t ds = int8_t(dspp <= 1 ? 1 : (dspp - 1) * (dspp - 1));

          adc_int -= (Phase2KinkADC - 1);
          adc_int *= ds;
          adc_int += (Phase2KinkADC - 1);

          adc_int = ((adc_int + 0.5 * ds) * ElectronPerADCGain);
        }

        adc_int += int(Phase2DigiBaseline);
      }

      adc[i] = std::min(adc_int, int(std::numeric_limits<uint16_t>::max()));
    }
  }

Variable Documentation

ElectronPerADCGain

constexpr float gpuCalibPixel::ElectronPerADCGain = 600

Definition at line 26 of file gpuCalibPixel.h.

Referenced by calibDigisPhase2().

Phase2DigiBaseline

constexpr uint16_t gpuCalibPixel::Phase2DigiBaseline = 1500

Definition at line 28 of file gpuCalibPixel.h.

Referenced by calibDigisPhase2().

Phase2KinkADC

constexpr uint8_t gpuCalibPixel::Phase2KinkADC = 8

Definition at line 29 of file gpuCalibPixel.h.

Referenced by calibDigisPhase2().

Phase2ReadoutMode

constexpr int8_t gpuCalibPixel::Phase2ReadoutMode = 3

Definition at line 27 of file gpuCalibPixel.h.

Referenced by calibDigisPhase2().

VCalChargeThreshold

constexpr int gpuCalibPixel::VCalChargeThreshold = 100

Definition at line 24 of file gpuCalibPixel.h.

VCaltoElectronGain

constexpr float gpuCalibPixel::VCaltoElectronGain = 47

Definition at line 20 of file gpuCalibPixel.h.

Referenced by calibDigis().

VCaltoElectronGain_L1

constexpr float gpuCalibPixel::VCaltoElectronGain_L1 = 50

Definition at line 21 of file gpuCalibPixel.h.

Referenced by calibDigis().

VCaltoElectronOffset

constexpr float gpuCalibPixel::VCaltoElectronOffset = -60

Definition at line 22 of file gpuCalibPixel.h.

Referenced by calibDigis().

VCaltoElectronOffset_L1

constexpr float gpuCalibPixel::VCaltoElectronOffset_L1 = -670

Definition at line 23 of file gpuCalibPixel.h.

Referenced by calibDigis().