diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc index 95aac36dbd197..993840c62c7f1 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc @@ -71,6 +71,7 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer wordFedAppender_; PixelDataFormatter::Errors errors_; + const bool isRun2_; const bool includeErrors_; const bool useQuality_; const bool usePilotBlade_; @@ -84,6 +85,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi gainsToken_(esConsumes()), cablingMapToken_(esConsumes( edm::ESInputTag("", iConfig.getParameter("CablingMapLabel")))), + isRun2_(iConfig.getParameter("isRun2")), includeErrors_(iConfig.getParameter("IncludeErrors")), useQuality_(iConfig.getParameter("UseQualityInfo")), usePilotBlade_(iConfig.getParameter("UsePilotBlade")) // Control the usage of pilot-blade data, FED=40 @@ -108,6 +110,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; + desc.add("isRun2", true); desc.add("IncludeErrors", true); desc.add("UseQualityInfo", false); desc.add("UsePilotBlade", false)->setComment("## Use pilot blades"); @@ -233,7 +236,8 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent, } // end of for loop - gpuAlgo_.makeClustersAsync(gpuMap, + gpuAlgo_.makeClustersAsync(isRun2_, + gpuMap, gpuModulesToUnpack, gpuGains, *wordFedAppender_, diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index cf2b10b198692..f14808dda1e2b 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -523,7 +523,8 @@ namespace pixelgpudetails { } // Interface to outside - void SiPixelRawToClusterGPUKernel::makeClustersAsync(const SiPixelFedCablingMapGPU *cablingMap, + void SiPixelRawToClusterGPUKernel::makeClustersAsync(bool isRun2, + const SiPixelFedCablingMapGPU *cablingMap, const unsigned char *modToUnp, const SiPixelGainForHLTonGPU *gains, const WordFedAppender &wordFed, @@ -599,7 +600,8 @@ namespace pixelgpudetails { int blocks = (std::max(int(wordCounter), int(gpuClustering::MaxNumModules)) + threadsPerBlock - 1) / threadsPerBlock; - gpuCalibPixel::calibDigis<<>>(digis_d.moduleInd(), + gpuCalibPixel::calibDigis<<>>(isRun2, + digis_d.moduleInd(), digis_d.c_xx(), digis_d.c_yy(), digis_d.adc(), diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index ee9729f75aed2..d214e7784af48 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h @@ -168,7 +168,8 @@ namespace pixelgpudetails { SiPixelRawToClusterGPUKernel& operator=(const SiPixelRawToClusterGPUKernel&) = delete; SiPixelRawToClusterGPUKernel& operator=(SiPixelRawToClusterGPUKernel&&) = delete; - void makeClustersAsync(const SiPixelFedCablingMapGPU* cablingMap, + void makeClustersAsync(bool isRun2, + const SiPixelFedCablingMapGPU* cablingMap, const unsigned char* modToUnp, const SiPixelGainForHLTonGPU* gains, const WordFedAppender& wordFed, diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h index 41e028b3c4595..50c62f44f1df8 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h @@ -13,12 +13,14 @@ namespace gpuCalibPixel { constexpr uint16_t InvId = 9999; // must be > MaxNumModules + // valid for run2 constexpr float VCaltoElectronGain = 47; // L2-4: 47 +- 4.7 constexpr float VCaltoElectronGain_L1 = 50; // L1: 49.6 +- 2.6 constexpr float VCaltoElectronOffset = -60; // L2-4: -60 +- 130 constexpr float VCaltoElectronOffset_L1 = -670; // L1: -670 +- 220 - __global__ void calibDigis(uint16_t* id, + __global__ void calibDigis(bool isRun2, + uint16_t* id, uint16_t const* __restrict__ x, uint16_t const* __restrict__ y, uint16_t* adc, @@ -41,8 +43,8 @@ namespace gpuCalibPixel { if (InvId == id[i]) continue; - float conversionFactor = id[i] < 96 ? VCaltoElectronGain_L1 : VCaltoElectronGain; - float offset = id[i] < 96 ? VCaltoElectronOffset_L1 : VCaltoElectronOffset; + 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; diff --git a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py index c80f3b16b3a43..3f8cf314ec2e2 100644 --- a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py +++ b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py @@ -9,6 +9,11 @@ siPixelClustersPreSplittingTask = cms.Task(siPixelClustersPreSplitting) siPixelClustersCUDAPreSplitting = _siPixelRawToClusterCUDA.clone() +from Configuration.Eras.Modifier_run3_common_cff import run3_common +run3_common.toModify(siPixelClustersCUDAPreSplitting, + isRun2=False +) + siPixelDigisClustersPreSplitting = _siPixelDigisClustersFromSoA.clone() siPixelClustersPreSplittingTaskCUDA = cms.Task( siPixelClustersCUDAPreSplitting,