Update Pixel gain calibration scheme (for Run3) (#492)

Update the Patatrack code following cms-sw#29333: Modify the scheme of the pixel gain calibration: instead of applying the VCal calibration in the clusterizer include it already in the gain calibration payload.
cms-patatrack · Dec 29, 2020 · 6ad3fec · 6ad3fec
1 parent e9ed9fa
commit 6ad3fec
Show file tree

Hide file tree

Showing 5 changed files with 21 additions and 7 deletions.
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
@@ -71,6 +71,7 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
   std::unique_ptr<pixelgpudetails::SiPixelRawToClusterGPUKernel::WordFedAppender> 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<SiPixelGainCalibrationForHLTGPU, SiPixelGainCalibrationForHLTGPURcd>()),
       cablingMapToken_(esConsumes<SiPixelFedCablingMap, SiPixelFedCablingMapRcd>(
           edm::ESInputTag("", iConfig.getParameter<std::string>("CablingMapLabel")))),
+      isRun2_(iConfig.getParameter<bool>("isRun2")),
       includeErrors_(iConfig.getParameter<bool>("IncludeErrors")),
       useQuality_(iConfig.getParameter<bool>("UseQualityInfo")),
       usePilotBlade_(iConfig.getParameter<bool>("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<bool>("isRun2", true);
   desc.add<bool>("IncludeErrors", true);
   desc.add<bool>("UseQualityInfo", false);
   desc.add<bool>("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
@@ -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<<<blocks, threadsPerBlock, 0, stream>>>(digis_d.moduleInd(),
+      gpuCalibPixel::calibDigis<<<blocks, threadsPerBlock, 0, stream>>>(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
@@ -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
@@ -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
@@ -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,