Merge pull request #33259 from czangela/pixel_local_reco_refactorings…

…_issue2_22_03_2021 GPU pixel local reconstruction: make the pixel cluster thresholds configurable
cms-sw · Apr 8, 2021 · eae9f05 · eae9f05
2 parents 55c4efe + 822c264
commit eae9f05
Show file tree

Hide file tree

Showing 7 changed files with 62 additions and 10 deletions.
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h
@@ -0,0 +1,14 @@
+#ifndef RecoLocalTracker_SiPixelClusterizer_plugins_SiPixelClusterThresholds_h
+#define RecoLocalTracker_SiPixelClusterizer_plugins_SiPixelClusterThresholds_h
+
+struct SiPixelClusterThresholds {
+  inline constexpr int32_t getThresholdForLayerOnCondition(bool isLayer1) const noexcept {
+    return isLayer1 ? layer1 : otherLayers;
+  }
+  const int32_t layer1;
+  const int32_t otherLayers;
+};
+
+constexpr SiPixelClusterThresholds kSiPixelClusterThresholdsDefaultPhase1{.layer1 = 2000, .otherLayers = 4000};
+
+#endif  // RecoLocalTracker_SiPixelClusterizer_plugins_SiPixelClusterThresholds_h
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc
@@ -15,7 +15,10 @@
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
 #include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
 #include "Geometry/Records/interface/TrackerTopologyRcd.h"
-#include "RecoLocalTracker/SiPixelClusterizer/plugins/PixelClusterizerBase.h"
+
+// local include(s)
+#include "PixelClusterizerBase.h"
+#include "SiPixelClusterThresholds.h"
 
 class SiPixelDigisClustersFromSoA : public edm::global::EDProducer<> {
 public:
@@ -33,17 +36,23 @@ class SiPixelDigisClustersFromSoA : public edm::global::EDProducer<> {
 
   edm::EDPutTokenT<edm::DetSetVector<PixelDigi>> digiPutToken_;
   edm::EDPutTokenT<SiPixelClusterCollectionNew> clusterPutToken_;
+
+  const SiPixelClusterThresholds clusterThresholds_;  // Cluster threshold in electrons
 };
 
 SiPixelDigisClustersFromSoA::SiPixelDigisClustersFromSoA(const edm::ParameterSet& iConfig)
     : topoToken_(esConsumes()),
       digiGetToken_(consumes<SiPixelDigisSoA>(iConfig.getParameter<edm::InputTag>("src"))),
       digiPutToken_(produces<edm::DetSetVector<PixelDigi>>()),
-      clusterPutToken_(produces<SiPixelClusterCollectionNew>()) {}
+      clusterPutToken_(produces<SiPixelClusterCollectionNew>()),
+      clusterThresholds_{iConfig.getParameter<int>("clusterThreshold_layer1"),
+                         iConfig.getParameter<int>("clusterThreshold_otherLayers")} {}
 
 void SiPixelDigisClustersFromSoA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
   edm::ParameterSetDescription desc;
   desc.add<edm::InputTag>("src", edm::InputTag("siPixelDigisSoA"));
+  desc.add<int>("clusterThreshold_layer1", kSiPixelClusterThresholdsDefaultPhase1.layer1);
+  desc.add<int>("clusterThreshold_otherLayers", kSiPixelClusterThresholdsDefaultPhase1.otherLayers);
   descriptions.addWithDefaultLabel(desc);
 }
 
@@ -77,7 +86,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con
       return;  // this in reality should never happen
     edmNew::DetSetVector<SiPixelCluster>::FastFiller spc(*outputClusters, detId);
     auto layer = (DetId(detId).subdetId() == 1) ? ttopo.pxbLayer(detId) : 0;
-    auto clusterThreshold = (layer == 1) ? 2000 : 4000;
+    auto clusterThreshold = clusterThresholds_.getThresholdForLayerOnCondition(layer == 1);
     for (int32_t ic = 0; ic < nclus + 1; ++ic) {
       auto const& acluster = aclusters[ic];
       // in any case we cannot  go out of sync with gpu...

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
@@ -37,6 +37,7 @@
 #include "RecoTracker/Record/interface/CkfComponentsRecord.h"
 
 // local includes
+#include "SiPixelClusterThresholds.h"
 #include "SiPixelRawToClusterGPUKernel.h"
 
 class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork> {
@@ -77,6 +78,7 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
   const bool isRun2_;
   const bool includeErrors_;
   const bool useQuality_;
+  const SiPixelClusterThresholds clusterThresholds_;
 };
 
 SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig)
@@ -89,7 +91,9 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi
           edm::ESInputTag("", iConfig.getParameter<std::string>("CablingMapLabel")))),
       isRun2_(iConfig.getParameter<bool>("isRun2")),
       includeErrors_(iConfig.getParameter<bool>("IncludeErrors")),
-      useQuality_(iConfig.getParameter<bool>("UseQualityInfo")) {
+      useQuality_(iConfig.getParameter<bool>("UseQualityInfo")),
+      clusterThresholds_{iConfig.getParameter<int32_t>("clusterThreshold_layer1"),
+                         iConfig.getParameter<int32_t>("clusterThreshold_otherLayers")} {
   if (includeErrors_) {
     digiErrorPutToken_ = produces<cms::cuda::Product<SiPixelDigiErrorsCUDA>>();
   }
@@ -110,6 +114,8 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d
   desc.add<bool>("isRun2", true);
   desc.add<bool>("IncludeErrors", true);
   desc.add<bool>("UseQualityInfo", false);
+  desc.add<int32_t>("clusterThreshold_layer1", kSiPixelClusterThresholdsDefaultPhase1.layer1);
+  desc.add<int32_t>("clusterThreshold_otherLayers", kSiPixelClusterThresholdsDefaultPhase1.otherLayers);
   desc.add<edm::InputTag>("InputLabel", edm::InputTag("rawDataCollector"));
   {
     edm::ParameterSetDescription psd0;
@@ -231,6 +237,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent,
   }  // end of for loop
 
   gpuAlgo_.makeClustersAsync(isRun2_,
+                             clusterThresholds_,
                              gpuMap,
                              gpuModulesToUnpack,
                              gpuGains,

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
@@ -510,6 +510,7 @@ namespace pixelgpudetails {
 
   // Interface to outside
   void SiPixelRawToClusterGPUKernel::makeClustersAsync(bool isRun2,
+                                                       const SiPixelClusterThresholds clusterThresholds,
                                                        const SiPixelROCsStatusAndMapping *cablingMap,
                                                        const unsigned char *modToUnp,
                                                        const SiPixelGainForHLTonGPU *gains,
@@ -635,7 +636,8 @@ namespace pixelgpudetails {
 #endif
 
       // apply charge cut
-      clusterChargeCut<<<blocks, threadsPerBlock, 0, stream>>>(digis_d.moduleInd(),
+      clusterChargeCut<<<blocks, threadsPerBlock, 0, stream>>>(clusterThresholds,
+                                                               digis_d.moduleInd(),
                                                                digis_d.adc(),
                                                                clusters_d.moduleStart(),
                                                                clusters_d.clusInModule(),

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
@@ -14,6 +14,9 @@
 #include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h"
 #include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h"
 
+// local include(s)
+#include "SiPixelClusterThresholds.h"
+
 struct SiPixelROCsStatusAndMapping;
 class SiPixelGainForHLTonGPU;
 
@@ -170,6 +173,7 @@ namespace pixelgpudetails {
     SiPixelRawToClusterGPUKernel& operator=(SiPixelRawToClusterGPUKernel&&) = delete;
 
     void makeClustersAsync(bool isRun2,
+                           const SiPixelClusterThresholds clusterThresholds,
                            const SiPixelROCsStatusAndMapping* cablingMap,
                            const unsigned char* modToUnp,
                            const SiPixelGainForHLTonGPU* gains,

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h
@@ -5,14 +5,20 @@
 #include <cstdio>
 
 #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
+#include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/prefixScan.h"
 
+// local include(s)
+#include "SiPixelClusterThresholds.h"
+
 namespace gpuClustering {
 
   __global__ void clusterChargeCut(
-      uint16_t* __restrict__ id,                 // module id of each pixel (modified if bad cluster)
-      uint16_t const* __restrict__ adc,          //  charge of each pixel
+      SiPixelClusterThresholds
+          clusterThresholds,             // charge cut on cluster in electrons (for layer 1 and for other layers)
+      uint16_t* __restrict__ id,         // module id of each pixel (modified if bad cluster)
+      uint16_t const* __restrict__ adc,  //  charge of each pixel
       uint32_t const* __restrict__ moduleStart,  // index of the first pixel of each module
       uint32_t* __restrict__ nClustersInModule,  // modified: number of clusters found in each module
       uint32_t const* __restrict__ moduleId,     // module id of each module
@@ -79,7 +85,8 @@ namespace gpuClustering {
       }
       __syncthreads();
 
-      auto chargeCut = thisModuleId < 96 ? 2000 : 4000;  // move in constants (calib?)
+      auto chargeCut =
+          clusterThresholds.getThresholdForLayerOnCondition(thisModuleId < phase1PixelTopology::layerStart[1]);
       for (auto i = threadIdx.x; i < nclus; i += blockDim.x) {
         newclusId[i] = ok[i] = charge[i] > chargeCut ? 1 : 0;
       }

diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h
@@ -18,6 +18,7 @@
 
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h"
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h"
+#include "RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h"
 
 int main(void) {
 #ifdef __CUDACC__
@@ -27,6 +28,7 @@ int main(void) {
   using namespace gpuClustering;
 
   constexpr int numElements = 256 * maxNumModules;
+  constexpr SiPixelClusterThresholds clusterThresholds(kSiPixelClusterThresholdsDefaultPhase1);
 
   // these in reality are already on GPU
   auto h_id = std::make_unique<uint16_t[]>(numElements);
@@ -288,6 +290,7 @@ int main(void) {
 
     cms::cuda::launch(clusterChargeCut,
                       {blocksPerGrid, threadsPerBlock},
+                      clusterThresholds,
                       d_id.get(),
                       d_adc.get(),
                       d_moduleStart.get(),
@@ -317,8 +320,14 @@ int main(void) {
     if (ncl != std::accumulate(nclus, nclus + maxNumModules, 0))
       std::cout << "ERROR!!!!! wrong number of cluster found" << std::endl;
 
-    clusterChargeCut(
-        h_id.get(), h_adc.get(), h_moduleStart.get(), h_clusInModule.get(), h_moduleId.get(), h_clus.get(), n);
+    clusterChargeCut(clusterThresholds,
+                     h_id.get(),
+                     h_adc.get(),
+                     h_moduleStart.get(),
+                     h_clusInModule.get(),
+                     h_moduleId.get(),
+                     h_clus.get(),
+                     n);
 #endif  // __CUDACC__
 
     std::cout << "found " << nModules << " Modules active" << std::endl;