Migrated PixelRecHit to Heterogeneous producer (#81)

Migrate PixelRecHit EDProducer to HeterogeneousEDProducer, including the cpu product. Data structures on gpu now include everything needed for Doublets, CA and fit. Layer splitting done: phi sorting (or partial sorting) requires #69. Includes some cleanup and bug fixes.
cms-patatrack · Dec 29, 2020 · 7501972 · 7501972
1 parent 4117157
commit 7501972
Show file tree

Hide file tree

Showing 4 changed files with 48 additions and 18 deletions.
diff --git a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h
@@ -20,6 +20,15 @@ namespace phase1PixelTopology {
 
   constexpr uint32_t numPixsInModule = uint32_t(numRowsInModule)* uint32_t(numColsInModule);
 
+  constexpr uint32_t numberOfModules = 1856;
+
+  constexpr uint32_t layerStart[11] = {0,96,320,672,1184,1296,1408,1520,1632,1744,1856};
+  constexpr char const * layerName[10] = {"BL1","BL2","BL3","BL4",
+   		  	                  "E+1", "E+2", "E+3",
+			                  "E-1", "E-2", "E-3"
+                                          };
+
+
   // this is for the ROC n<512 (upgrade 1024)
   constexpr inline
   uint16_t  divu52(uint16_t n) {

diff --git a/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py b/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py
@@ -17,16 +17,23 @@
 striptrackerlocalreco = cms.Sequence(siStripZeroSuppression*siStripClusters*siStripMatchedRecHits)
 trackerlocalreco = cms.Sequence(pixeltrackerlocalreco*striptrackerlocalreco*clusterSummaryProducer)
 
+
 from RecoLocalTracker.SiPixelClusterizer.siPixelClustersHeterogeneous_cfi import *
 from RecoLocalTracker.SiPixelClusterizer.siPixelFedCablingMapGPUWrapper_cfi import *
 from CalibTracker.SiPixelESProducers.siPixelGainCalibrationForHLTGPU_cfi import *
 
+from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import *
+from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneousConverter_cfi import siPixelRecHitHeterogeneousConverter as _siPixelRecHitHeterogeneousConverter
+gpu.toReplaceWith(siPixelRecHitsPreSplitting, _siPixelRecHitHeterogeneousConverter.clone())
+
+
+
 from Configuration.ProcessModifiers.gpu_cff import gpu
 _pixeltrackerlocalreco_gpu = pixeltrackerlocalreco.copy()
 _pixeltrackerlocalreco_gpu.replace(siPixelClustersPreSplitting, siPixelClustersHeterogeneous+siPixelClustersPreSplitting)
+_pixeltrackerlocalreco_gpu.replace(siPixelRecHitsPreSplitting, siPixelRecHitHeterogeneous+siPixelRecHitsPreSplitting)
 gpu.toReplaceWith(pixeltrackerlocalreco, _pixeltrackerlocalreco_gpu)
 
-
 from RecoLocalTracker.SiPhase2Clusterizer.phase2TrackerClusterizer_cfi import *
 from RecoLocalTracker.Phase2TrackerRecHits.Phase2StripCPEGeometricESProducer_cfi import *
 

diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
@@ -6,10 +6,14 @@
 #include <cstdio>
 #include <limits>
 
+#include "DataFormats/Math/interface/approx_atan2.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
 
 namespace gpuPixelRecHits {
 
+
+
+
   // to be moved in common namespace...
   constexpr uint16_t InvId=9999; // must be > MaxNumModules
 
@@ -20,20 +24,23 @@ namespace gpuPixelRecHits {
 
 
   __global__ void getHits(pixelCPEforGPU::ParamsOnGPU const * cpeParams,
+                          float const * bs,
                           uint16_t const * id,
 			  uint16_t const * x,
 			  uint16_t const * y,
 			  uint16_t const * adc,
 			  uint32_t const * digiModuleStart,
 			  uint32_t const * clusInModule,
 			  uint32_t const * moduleId,
-			  int32_t const * clus,
+			  int32_t  const * clus,
 			  int numElements,
 			  uint32_t const * hitsModuleStart,
                           int32_t * chargeh,
-			  float * xh, float * yh, float * zh,
-                          float * xe, float * ye, uint16_t * mr,
-			  bool local)   // if true fill just x & y in local coord
+                          uint16_t * detInd,
+			  float * xg, float * yg, float * zg, float * rg, int16_t * iph,
+                          float * xl, float * yl,
+                          float * xe, float * ye, 
+                          uint16_t * mr, uint16_t * mc)
   {
     // as usual one block per module
     __shared__ ClusParams clusParams;
@@ -108,16 +115,26 @@ namespace gpuPixelRecHits {
 
     chargeh[h] = clusParams.charge[ic];
 
-    if (local) {
-      xh[h] = clusParams.xpos[ic];
-      yh[h] = clusParams.ypos[ic];
-    } else {
-      cpeParams->detParams(me).frame.toGlobal(clusParams.xpos[ic], clusParams.ypos[ic],
-                                              xh[h], yh[h], zh[h] );
-    }
-    xe[h] = clusParams.xerr[ic];
-    ye[h] = clusParams.yerr[ic];
-    mr[h] = clusParams.minRow[ic];
+    detInd[h] = me;
+
+    xl[h]= clusParams.xpos[ic];   
+    yl[h]= clusParams.ypos[ic]; 
+
+    xe[h]= clusParams.xerr[ic];
+    ye[h]= clusParams.yerr[ic];
+    mr[h]= clusParams.minRow[ic];
+    mc[h]= clusParams.minCol[ic];
+
+    // to global and compute phi... 
+    cpeParams->detParams(me).frame.toGlobal(xl[h],yl[h], xg[h],yg[h],zg[h]);
+    // here correct for the beamspot...
+    xg[h]-=bs[0];
+    yg[h]-=bs[1];
+    zg[h]-=bs[2];
+
+    rg[h] = std::sqrt(xg[h]*xg[h]+yg[h]*yg[h]);
+    iph[h] = unsafe_atan2s<7>(yg[h],xg[h]);
+
   }
 
 }

diff --git a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py
@@ -7,9 +7,6 @@
     VerboseLevel = cms.untracked.int32(0),
 )
 
-from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import siPixelRecHitHeterogeneous as _siPixelRecHitHeterogeneous
-gpu.toReplaceWith(siPixelRecHits, _siPixelRecHitHeterogeneous)
-
 siPixelRecHitsPreSplitting = siPixelRecHits.clone(
     src = 'siPixelClustersPreSplitting'
 )