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

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 30 of file gpuCalibPixel.h.

References cms::cudacompat::blockDim, cms::cudacompat::blockIdx, cuy::col, first, cms::cudacompat::gridDim, mps_fire::i, gpuClustering::invalidModuleId, HLT_FULL_cff::isRun2, SiStripPI::max, phase1PixelTopology::numberOfModules, gpuClustering::numElements, hltrates_dqm_sourceclient-live_cfg::offset, EcalCondDBWriter_cfi::pedestal, gpuVertexFinder::printf(), runTheMatrix::ret, cms::cudacompat::threadIdx, VCaltoElectronGain, VCaltoElectronGain_L1, VCaltoElectronOffset, and VCaltoElectronOffset_L1.

     {
     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::max(100, int(vcal));
       }
     }
   }

__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 77 of file gpuCalibPixel.h.

References cms::cudacompat::blockDim, cms::cudacompat::blockIdx, ElectronPerADCGain, first, cms::cudacompat::gridDim, mps_fire::i, gpuClustering::invalidModuleId, universalConfigTemplate::mode, 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);
 
       if constexpr (mode < 0)
         adc[i] = int(adc[i] * ElectronPerADCGain);
       else {
         if (adc[i] < Phase2KinkADC)
           adc[i] = int((adc[i] - 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[i] -= (Phase2KinkADC - 1);
           adc[i] *= ds;
           adc[i] += (Phase2KinkADC - 1);
 
           adc[i] = uint16_t((adc[i] + 0.5 * ds) * ElectronPerADCGain);
         }
 
         adc[i] += int(Phase2DigiBaseline);
       }
     }
   }