gpuCalibPixel Namespace Reference

Functions
__global__ void	calibDigis (bool isRun2, 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)

Variables
constexpr float	VCaltoElectronGain = 47
constexpr float	VCaltoElectronGain_L1 = 50
constexpr float	VCaltoElectronOffset = -60
constexpr float	VCaltoElectronOffset_L1 = -670

Function Documentation

calibDigis()

__global__ void gpuCalibPixel::calibDigis	(	bool	isRun2,
		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 21 of file gpuCalibPixel.h.

    {
    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 < gpuClustering::maxNumModules; i += gridDim.x * blockDim.x) {
      nClustersInModule[i] = 0;
    }
 
    for (int i = first; i < numElements; i += gridDim.x * blockDim.x) {
      if (invalidModuleId == id[i])
        continue;
 
      float conversionFactor = (isRun2) ? (id[i] < 96 ? VCaltoElectronGain_L1 : VCaltoElectronGain) : 1.f;
      float offset = (isRun2) ? (id[i] < 96 ? VCaltoElectronOffset_L1 : VCaltoElectronOffset) : 0;
 
      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) {
        id[i] = invalidModuleId;
        adc[i] = 0;
        printf("bad pixel at %d in %d\n", i, id[i]);
      } else {
        float vcal = adc[i] * gain - pedestal * gain;
        adc[i] = std::max(100, int(vcal * conversionFactor + offset));
      }
    }
  }

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

Variable Documentation

VCaltoElectronGain

constexpr float gpuCalibPixel::VCaltoElectronGain = 47

constexpr

Definition at line 16 of file gpuCalibPixel.h.

Referenced by calibDigis().

VCaltoElectronGain_L1

constexpr float gpuCalibPixel::VCaltoElectronGain_L1 = 50

constexpr

Definition at line 17 of file gpuCalibPixel.h.

Referenced by calibDigis().

VCaltoElectronOffset

constexpr float gpuCalibPixel::VCaltoElectronOffset = -60

constexpr

Definition at line 18 of file gpuCalibPixel.h.

Referenced by calibDigis().

VCaltoElectronOffset_L1

constexpr float gpuCalibPixel::VCaltoElectronOffset_L1 = -670

constexpr

Definition at line 19 of file gpuCalibPixel.h.

Referenced by calibDigis().