Add truncation to pixel charge on GPU (#501)

Makes the GPU implementation of CPE and CPEFast consistent with the CPEGeneric one.
cms-patatrack · Dec 29, 2020 · 1ab9b2d · 1ab9b2d
1 parent c279c8a
commit 1ab9b2d
Show file tree

Hide file tree

Showing 3 changed files with 11 additions and 1 deletion.
diff --git a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h
@@ -36,6 +36,7 @@ namespace pixelCPEforGPU {
     float shiftY;
     float chargeWidthX;
     float chargeWidthY;
+    uint16_t pixmx;  // max pix charge
 
     float x0, y0, z0;  // the vertex in the local coord of the detector
 

diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
@@ -134,6 +134,7 @@ namespace gpuPixelRecHits {
 
       __syncthreads();
 
+      auto pixmx = cpeParams->detParams(me).pixmx;
       for (int i = first; i < numElements; i += blockDim.x) {
         auto id = digis.moduleInd(i);
         if (id == InvId)
@@ -148,7 +149,7 @@ namespace gpuPixelRecHits {
         assert(cl < MaxHitsInIter);
         auto x = digis.xx(i);
         auto y = digis.yy(i);
-        auto ch = digis.adc(i);
+        auto ch = std::min(digis.adc(i), pixmx);
         atomicAdd(&clusParams.charge[cl], ch);
         if (clusParams.minRow[cl] == x)
           atomicAdd(&clusParams.Q_f_X[cl], ch);

diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc
@@ -112,6 +112,8 @@ void PixelCPEFast::fillParamsForGpu() {
   m_commonParamsGPU.thePitchX = m_DetParams[0].thePitchX;
   m_commonParamsGPU.thePitchY = m_DetParams[0].thePitchY;
 
+  // std::cout << "pitch & thickness " <<  m_commonParamsGPU.thePitchX << ' ' << m_commonParamsGPU.thePitchY << "  " << m_commonParamsGPU.theThicknessB << ' ' << m_commonParamsGPU.theThicknessE << std::endl;
+
   // zero average geometry
   memset(&m_averageGeometry, 0, sizeof(pixelCPEforGPU::AverageGeometry));
 
@@ -210,6 +212,7 @@ void PixelCPEFast::fillParamsForGpu() {
 #endif
 
     errorFromTemplates(p, cp, 20000.f);
+    g.pixmx = std::max(0, cp.pixmx);
     g.sx[0] = cp.sigmax;
     g.sx[1] = cp.sx1;
     g.sx[2] = cp.sx2;
@@ -407,6 +410,9 @@ LocalPoint PixelCPEFast::localPosition(DetParam const& theDetParam, ClusterParam
   auto xPos = cp.xpos[0];
   auto yPos = cp.ypos[0];
 
+  //  std::cout<<" in PixelCPEFast:localPosition - pos = "<<xPos<<" "<<yPos
+  //           << " size "<< cp.maxRow[0]-cp.minRow[0] << ' ' << cp.maxCol[0]-cp.minCol[0] << std::endl; //dk
+
   //--- Now put the two together
   LocalPoint pos_in_local(xPos, yPos);
   return pos_in_local;
@@ -576,6 +582,8 @@ LocalError PixelCPEFast::localError(DetParam const& theDetParam, ClusterParam& t
 
   }  // end
 
+  //   std::cout<<" errors  "<<xerr<<" "<<yerr<<std::endl;  //dk
+
   auto xerr_sq = xerr * xerr;
   auto yerr_sq = yerr * yerr;