Make cluster charge cut thresholds configurable for the function gpuC…

…lusterChargeCut See cms-sw#32483
czangela · Mar 24, 2021 · b26dc4f · b26dc4f
1 parent 5b26e4a
commit b26dc4f
Show file tree

Hide file tree

Showing 6 changed files with 42 additions and 7 deletions.
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
@@ -35,6 +35,7 @@
 #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
 #include "RecoTracker/Record/interface/CkfComponentsRecord.h"
+#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h"
 
 // local includes
 #include "SiPixelRawToClusterGPUKernel.h"
@@ -77,6 +78,8 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
   const bool isRun2_;
   const bool includeErrors_;
   const bool useQuality_;
+
+  const gpuClusterThresholds clusterThresholds_;
 };
 
 SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig)
@@ -89,7 +92,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>("theGPUClusterThreshold_L1"),
+                         iConfig.getParameter<int32_t>("theGPUClusterThreshold")) {
   if (includeErrors_) {
     digiErrorPutToken_ = produces<cms::cuda::Product<SiPixelDigiErrorsCUDA>>();
   }
@@ -110,6 +115,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>("theGPUClusterThreshold_L1", 2000);
+  desc.add<int32_t>("theGPUClusterThreshold", 4000);
   desc.add<edm::InputTag>("InputLabel", edm::InputTag("rawDataCollector"));
   {
     edm::ParameterSetDescription psd0;
@@ -231,6 +238,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 gpuClusterThresholds 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
@@ -13,6 +13,7 @@
 #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h"
 #include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h"
 #include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h"
+#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h"
 
 struct SiPixelROCsStatusAndMapping;
 class SiPixelGainForHLTonGPU;
@@ -170,6 +171,7 @@ namespace pixelgpudetails {
     SiPixelRawToClusterGPUKernel& operator=(SiPixelRawToClusterGPUKernel&&) = delete;
 
     void makeClustersAsync(bool isRun2,
+                           const gpuClusterThresholds 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
@@ -7,12 +7,14 @@
 #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/prefixScan.h"
+#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.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
+      gpuClusterThresholds 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 +81,7 @@ namespace gpuClustering {
       }
       __syncthreads();
 
-      auto chargeCut = thisModuleId < 96 ? 2000 : 4000;  // move in constants (calib?)
+      auto chargeCut = thisModuleId < 96 ? clusterThresholds.layer1 : clusterThresholds.otherLayers;
       for (auto i = threadIdx.x; i < nclus; i += blockDim.x) {
         newclusId[i] = ok[i] = charge[i] > chargeCut ? 1 : 0;
       }

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h
@@ -0,0 +1,11 @@
+#ifndef RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterThresholds_h
+#define RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterThresholds_h
+
+struct gpuClusterThresholds {
+    constexpr gpuClusterThresholds(int32_t clusterThreshold_L1, int32_t clusterThreshold):
+        layer1(clusterThreshold_L1), otherLayers(clusterThreshold){};
+    const int32_t layer1;           // Cluster threshold in electrons for Layer 1
+    const int32_t otherLayers;      // Cluster threshold in electrons
+};
+
+#endif // RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterThresholds_h
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/gpuClusterThresholds.h"
 
 int main(void) {
 #ifdef __CUDACC__
@@ -28,6 +29,8 @@ int main(void) {
 
   constexpr int numElements = 256 * maxNumModules;
 
+  constexpr gpuClusterThresholds clusterThresholds(2000, 4000);
+
   // these in reality are already on GPU
   auto h_id = std::make_unique<uint16_t[]>(numElements);
   auto h_x = std::make_unique<uint16_t[]>(numElements);
@@ -288,6 +291,7 @@ int main(void) {
 
     cms::cuda::launch(clusterChargeCut,
                       {blocksPerGrid, threadsPerBlock},
+                      clusterThresholds,
                       d_id.get(),
                       d_adc.get(),
                       d_moduleStart.get(),
@@ -317,8 +321,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;