From 27bbcdd746ef2bc3a1833130c674ba80f7990f4c Mon Sep 17 00:00:00 2001
From: Vincenzo Innocente <vincenzo.innocente@cern.ch>
Date: Tue, 8 Jan 2019 18:29:08 +0100
Subject: [PATCH] Full workflow from raw data to pixel tracks and vertices on
 GPUs (cms-patatrack#216)

Port and optimise the full workflow from pixel raw data to pixel tracks and vertices to GPUs.
Clean the pixel n-tuplets with the "fishbone" algorithm (only on GPUs).

Other changes:
  - recover the Riemann fit updates lost during the merge with CMSSW 10.4.x;
  - speed up clustering and track fitting;
  - minor bug fix to avoid trivial regression with the optimized fit.
---
 .../interface/phase1PixelTopology.h           |  61 +++++
 .../test/phase1PixelTopology_t.cpp            | 231 +++++++++---------
 .../plugins/gpuClustering.h                   |  46 +++-
 .../SiPixelRecHits/interface/PixelCPEBase.h   |   5 +-
 .../SiPixelRecHits/interface/PixelCPEFast.h   |   4 +
 .../SiPixelRecHits/interface/pixelCPEforGPU.h |  44 +++-
 .../SiPixelRecHits/plugins/gpuPixelRecHits.h  |   6 +-
 .../SiPixelRecHits/src/PixelCPEFast.cc        | 130 ++++++++--
 8 files changed, 366 insertions(+), 161 deletions(-)

diff --git a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h
index 37c97a92a3eaa..68fb60361d40d 100644
--- a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h
+++ b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h
@@ -2,6 +2,7 @@
 #define Geometry_TrackerGeometryBuilder_phase1PixelTopology_h
 
 #include <cstdint>
+#include <array>
 
 namespace phase1PixelTopology {
 
@@ -29,6 +30,66 @@ namespace phase1PixelTopology {
                                           };
 
 
+  template<class Function, std::size_t... Indices>
+  constexpr auto map_to_array_helper(Function f, std::index_sequence<Indices...>)
+  -> std::array<typename std::result_of<Function(std::size_t)>::type, sizeof...(Indices)>
+  {
+    return {{ f(Indices)... }};
+  }
+
+  template<int N, class Function>
+  constexpr auto map_to_array(Function f)
+  -> std::array<typename std::result_of<Function(std::size_t)>::type, N>
+  {
+    return map_to_array_helper(f, std::make_index_sequence<N>{});
+  }
+
+
+  constexpr uint32_t findMaxModuleStride() {
+    bool go = true;
+    int n=2;
+    while (go) {
+      for  (uint8_t i=1; i<11; ++i) {
+        if (layerStart[i]%n !=0) {go=false; break;}
+      }
+      if(!go) break;
+      n*=2;
+    }
+    return n/2;
+  }
+
+  constexpr uint32_t maxModuleStride = findMaxModuleStride();
+
+
+  constexpr uint8_t findLayer(uint32_t detId) {
+    for  (uint8_t i=0; i<11; ++i) if (detId<layerStart[i+1]) return i;
+    return 11;
+  }
+
+  constexpr uint8_t findLayerFromCompact(uint32_t detId) {
+    detId*=maxModuleStride;
+    for  (uint8_t i=0; i<11; ++i) if (detId<layerStart[i+1]) return i;
+    return 11;
+  }
+
+
+  constexpr uint32_t layerIndexSize = numberOfModules/maxModuleStride;
+  constexpr std::array<uint8_t,layerIndexSize> layer = map_to_array<layerIndexSize>(findLayerFromCompact);
+
+  constexpr bool validateLayerIndex() {
+    bool res=true;
+    for (auto i=0U; i<numberOfModules; ++i)  {
+      auto j = i/maxModuleStride;
+      res &=(layer[j]<10);
+      res &=(i>=layerStart[layer[j]]);
+      res &=(i<layerStart[layer[j]+1]);
+    }
+    return res;
+  }
+
+  static_assert(validateLayerIndex(),"layer from detIndex algo is buggy");
+
+ 
   // this is for the ROC n<512 (upgrade 1024)
   constexpr inline
   uint16_t  divu52(uint16_t n) {
diff --git a/Geometry/TrackerGeometryBuilder/test/phase1PixelTopology_t.cpp b/Geometry/TrackerGeometryBuilder/test/phase1PixelTopology_t.cpp
index 293febbbc7143..43128fd02e7f6 100644
--- a/Geometry/TrackerGeometryBuilder/test/phase1PixelTopology_t.cpp
+++ b/Geometry/TrackerGeometryBuilder/test/phase1PixelTopology_t.cpp
@@ -1,128 +1,128 @@
+#include <cassert>
+#include <iostream>
+#include <tuple>
+
 #include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h"
-#include<cassert>
-#include<tuple>
 
 namespace {
 
   // original code from CMSSW_4_4
 
   std::tuple<int,bool> localXori(int mpx) {
-   const float m_pitchx=1.f;
-   int binoffx = int(mpx);             // truncate to int
-   float local_pitchx = m_pitchx;      // defaultpitch
-
-   if (binoffx>80) {            // ROC 1 - handles x on edge cluster
-     binoffx=binoffx+2;
-   } else if (binoffx==80) {    // ROC 1
-     binoffx=binoffx+1;
-     local_pitchx = 2 * m_pitchx;
-     
-   } else if (binoffx==79) {      // ROC 0
-     binoffx=binoffx+0;
-     local_pitchx = 2 * m_pitchx;    
-   } else if (binoffx>=0) {       // ROC 0
-     binoffx=binoffx+0;
-     
-   } else { // too small
-       assert("binoffx too small"==0);
-   }
-
-   return std::make_tuple(binoffx,local_pitchx>m_pitchx);
-  }
+    const float m_pitchx=1.f;
+    int binoffx = int(mpx);             // truncate to int
+    float local_pitchx = m_pitchx;      // defaultpitch
+
+    if (binoffx>80) {                   // ROC 1 - handles x on edge cluster
+      binoffx=binoffx+2;
+    } else if (binoffx==80) {           // ROC 1
+      binoffx=binoffx+1;
+      local_pitchx = 2 * m_pitchx;
+
+    } else if (binoffx==79) {           // ROC 0
+      binoffx=binoffx+0;
+      local_pitchx = 2 * m_pitchx;
+    } else if (binoffx>=0) {            // ROC 0
+      binoffx=binoffx+0;
+
+    } else {                            // too small
+      assert("binoffx too small"==0);
+    }
 
+    return std::make_tuple(binoffx,local_pitchx>m_pitchx);
+  }
 
   std::tuple<int,bool> localYori(int mpy) {
-   const float m_pitchy=1.f;
-   int binoffy = int(mpy);             // truncate to int
-   float local_pitchy = m_pitchy;      // defaultpitch
- 
-   if (binoffy>416) {            // ROC 8, not real ROC
-     binoffy=binoffy+17;
-   } else if (binoffy==416) {    // ROC 8
-     binoffy=binoffy+16;
-     local_pitchy = 2 * m_pitchy;
-         
-   } else if (binoffy==415) {    // ROC 7, last big pixel
-     binoffy=binoffy+15;
-     local_pitchy = 2 * m_pitchy;
-   } else if (binoffy>364) {     // ROC 7
-     binoffy=binoffy+15;
-   } else if (binoffy==364) {    // ROC 7
-     binoffy=binoffy+14;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==363) {      // ROC 6
-     binoffy=binoffy+13;
-     local_pitchy = 2 * m_pitchy;    
-   } else if (binoffy>312) {       // ROC 6
-     binoffy=binoffy+13;
-   } else if (binoffy==312) {      // ROC 6
-     binoffy=binoffy+12;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==311) {      // ROC 5
-     binoffy=binoffy+11;
-     local_pitchy = 2 * m_pitchy;    
-   } else if (binoffy>260) {      // ROC 5
-     binoffy=binoffy+11;
-   } else if (binoffy==260) {     // ROC 5
-     binoffy=binoffy+10;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==259) {      // ROC 4
-     binoffy=binoffy+9;
-     local_pitchy = 2 * m_pitchy;    
-   } else if (binoffy>208) {       // ROC 4
-     binoffy=binoffy+9;
-   } else if (binoffy==208) {      // ROC 4
-     binoffy=binoffy+8;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==207)  {      // ROC 3
-     binoffy=binoffy+7;
-     local_pitchy = 2 * m_pitchy;    
-     } else if (binoffy>156) {       // ROC 3
-     binoffy=binoffy+7;
-   } else if (binoffy==156) {      // ROC 3
-     binoffy=binoffy+6;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==155) {      // ROC 2
-     binoffy=binoffy+5;
-     local_pitchy = 2 * m_pitchy;    
-   } else if (binoffy>104) {       // ROC 2
-     binoffy=binoffy+5;
-   } else if (binoffy==104) {      // ROC 2
-     binoffy=binoffy+4;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==103) {      // ROC 1
-     binoffy=binoffy+3;
-     local_pitchy = 2 * m_pitchy;    
-   } else if (binoffy>52) {       // ROC 1
-     binoffy=binoffy+3;
-   } else if (binoffy==52) {      // ROC 1
-     binoffy=binoffy+2;
-     local_pitchy = 2 * m_pitchy;
-     
-   } else if (binoffy==51) {      // ROC 0
-     binoffy=binoffy+1;
-     local_pitchy = 2 * m_pitchy;    
-   } else if (binoffy>0) {         // ROC 0
-     binoffy=binoffy+1;
-   } else if (binoffy==0) {       // ROC 0
-     binoffy=binoffy+0;
-     local_pitchy = 2 * m_pitchy;
-   } else {
-       assert("binoffy too small"==0);
-   } 
-
-   return std::make_tuple(binoffy,local_pitchy>m_pitchy);
-  }   
+    const float m_pitchy=1.f;
+    int binoffy = int(mpy);             // truncate to int
+    float local_pitchy = m_pitchy;      // defaultpitch
+
+    if (binoffy>416) {                  // ROC 8, not real ROC
+      binoffy=binoffy+17;
+    } else if (binoffy==416) {          // ROC 8
+      binoffy=binoffy+16;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==415) {          // ROC 7, last big pixel
+      binoffy=binoffy+15;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>364) {           // ROC 7
+      binoffy=binoffy+15;
+    } else if (binoffy==364) {          // ROC 7
+      binoffy=binoffy+14;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==363) {          // ROC 6
+      binoffy=binoffy+13;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>312) {           // ROC 6
+      binoffy=binoffy+13;
+    } else if (binoffy==312) {          // ROC 6
+      binoffy=binoffy+12;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==311) {          // ROC 5
+      binoffy=binoffy+11;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>260) {           // ROC 5
+      binoffy=binoffy+11;
+    } else if (binoffy==260) {          // ROC 5
+      binoffy=binoffy+10;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==259) {          // ROC 4
+      binoffy=binoffy+9;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>208) {           // ROC 4
+      binoffy=binoffy+9;
+    } else if (binoffy==208) {          // ROC 4
+      binoffy=binoffy+8;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==207)  {         // ROC 3
+      binoffy=binoffy+7;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>156) {           // ROC 3
+      binoffy=binoffy+7;
+    } else if (binoffy==156) {          // ROC 3
+      binoffy=binoffy+6;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==155) {          // ROC 2
+      binoffy=binoffy+5;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>104) {           // ROC 2
+      binoffy=binoffy+5;
+    } else if (binoffy==104) {          // ROC 2
+      binoffy=binoffy+4;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==103) {          // ROC 1
+      binoffy=binoffy+3;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>52) {            // ROC 1
+      binoffy=binoffy+3;
+    } else if (binoffy==52) {           // ROC 1
+      binoffy=binoffy+2;
+      local_pitchy = 2 * m_pitchy;
+
+    } else if (binoffy==51) {           // ROC 0
+      binoffy=binoffy+1;
+      local_pitchy = 2 * m_pitchy;
+    } else if (binoffy>0) {             // ROC 0
+      binoffy=binoffy+1;
+    } else if (binoffy==0) {            // ROC 0
+      binoffy=binoffy+0;
+      local_pitchy = 2 * m_pitchy;
+    } else {
+      assert("binoffy too small"==0);
+    }
+
+    return std::make_tuple(binoffy,local_pitchy>m_pitchy);
+  }
 
 }
 
-#include<iostream>
 int main() {
 
   for (uint16_t ix=0; ix<80*2; ++ix) {
@@ -141,6 +141,13 @@ int main() {
     assert(std::get<1>(ori)==bp);
   }
 
+  using namespace phase1PixelTopology;
+  for (auto i=0U; i<numberOfModules; ++i)  {
+    assert(layer[i]<10);
+    assert(i>=layerStart[layer[i]]);
+    assert(i<layerStart[layer[i]+1]);
+  }
+
 
   return 0;
 }
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h
index 26bb4f9244c6a..dba9fe12f5492 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h
@@ -81,10 +81,9 @@ namespace gpuClustering {
    constexpr uint32_t maxPixInModule = 4000;
    constexpr auto  nbins = phase1PixelTopology::numColsInModule + 2;   //2+2;
    using Hist = HistoContainer<uint16_t,nbins,maxPixInModule,9,uint16_t>;
-   constexpr auto wss = Hist::totbins();
     __shared__ Hist hist;
-    __shared__ typename Hist::Counter ws[wss];
-    for (auto j=threadIdx.x; j<Hist::totbins(); j+=blockDim.x) { hist.off[j]=0; ws[j]=0;}
+    __shared__ typename Hist::Counter ws[32];
+    for (auto j=threadIdx.x; j<Hist::totbins(); j+=blockDim.x) { hist.off[j]=0;}
     __syncthreads();
 
     assert((msize == numElements) or ((msize < numElements) and (id[msize] != thisModuleId)));
@@ -107,10 +106,10 @@ namespace gpuClustering {
 #endif
       }
     __syncthreads();
-    hist.finalize(ws);
-    __syncthreads();
     if (threadIdx.x<32) ws[threadIdx.x]=0;  // used by prefix scan...
     __syncthreads();
+    hist.finalize(ws);
+    __syncthreads();
 #ifdef GPU_DEBUG
     assert(hist.size()==totGood);
     if (thisModuleId % 100 == 1)
@@ -120,7 +119,7 @@ namespace gpuClustering {
       for (int i = first; i < msize; i += blockDim.x) {
         if (id[i] == InvId)                 // skip invalid pixels
           continue;
-        hist.fill(y[i],i-firstPixel,ws);
+        hist.fill(y[i],i-firstPixel);
       }
 
 #ifdef CLUS_LIMIT_LOOP
@@ -134,21 +133,45 @@ namespace gpuClustering {
       jmax[k] = hist.end();
 #endif
 
-#ifdef GPU_DEBUG
     __shared__ int nloops;
     nloops=0;
-#endif
 
 
     __syncthreads();  // for hit filling!
 
+#ifdef GPU_DEBUG
+    // look for anomalous high occupancy
+    __shared__ uint32_t n40,n60;
+    n40=n60=0;
+    __syncthreads();
+    for (auto j=threadIdx.x; j<Hist::nbins(); j+=blockDim.x) { 
+      if(hist.size(j)>60) atomicAdd(&n60,1);
+      if(hist.size(j)>40) atomicAdd(&n40,1);
+     }
+    __syncthreads();
+    if (0==threadIdx.x) {
+      if (n60>0) printf("columns with more than 60 px %d in %d\n",n60,thisModuleId);
+      else if (n40>0) printf("columns with more than 40 px %d in %d\n",n40,thisModuleId);
+    }
+    __syncthreads();
+#endif
+
     // for each pixel, look at all the pixels until the end of the module;
     // when two valid pixels within +/- 1 in x or y are found, set their id to the minimum;
     // after the loop, all the pixel in each cluster should have the id equeal to the lowest
     // pixel in the cluster ( clus[i] == i ).
     bool more = true;
     while (__syncthreads_or(more)) {
-      more = false;
+      if (1==nloops%2) {
+        for (int j=threadIdx.x, k = 0; j<hist.size(); j+=blockDim.x, ++k) {
+             auto p = hist.begin()+j;
+             auto i = *p + firstPixel;
+             auto m = clusterId[i];
+             while (m!=clusterId[m]) m=clusterId[m];
+             clusterId[i]=m;
+        }
+      } else {
+        more = false;
         for (int j=threadIdx.x, k = 0; j<hist.size(); j+=blockDim.x, ++k) {
           auto p = hist.begin()+j;
           auto i = *p + firstPixel;
@@ -166,9 +189,7 @@ namespace gpuClustering {
           // loop to columns
           auto loop = [&](uint16_t const * kk) {
             auto m = (*kk)+firstPixel;
-#ifdef GPU_DEBUG
             assert(m!=i);
-#endif
             if (std::abs(int(x[m]) - int(x[i])) > 1) return;
             // if (std::abs(int(y[m]) - int(y[i])) > 1) return; // binssize is 1
             auto old = atomicMin(&clusterId[m], clusterId[i]);
@@ -185,9 +206,8 @@ namespace gpuClustering {
           ++p;
           for (;p<e;++p) loop(p);
         } // pixel loop
-#ifdef GPU_DEBUG
+        }
         if (threadIdx.x==0) ++nloops;
-#endif
     }  // end while
 
 #ifdef GPU_DEBUG
diff --git a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h
index 758880251f9f1..f908325029afe 100644
--- a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h
+++ b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h
@@ -80,11 +80,12 @@ class PixelCPEBase : public PixelClusterParameterEstimator
    
    struct ClusterParam
    {
+      ClusterParam(){}
       ClusterParam(const SiPixelCluster & cl) : theCluster(&cl) {}
 
       virtual ~ClusterParam() = default;
 
-      const SiPixelCluster * theCluster;
+      const SiPixelCluster * theCluster = nullptr;;
       
       //--- Cluster-level quantities (filled in computeAnglesFrom....)
       float cotalpha;
@@ -112,6 +113,7 @@ class PixelCPEBase : public PixelClusterParameterEstimator
       bool  hasFilledProb_ =false;
       // ggiurgiu@jhu.edu (10/18/2008)
       bool with_track_angle; // filled in computeAnglesFrom....
+      bool filled_from_2d = false; //
 
      // More detailed edge information (for CPE ClusterRepair, and elsewhere...)
      int   edgeTypeX_ = 0;   // 0: not on edge, 1: low end on edge, 2: high end
@@ -270,7 +272,6 @@ class PixelCPEBase : public PixelClusterParameterEstimator
    LocalVector driftDirection       (DetParam & theDetParam, LocalVector bfield ) const ;
    void computeLorentzShifts(DetParam &) const ;
    
-   bool isFlipped(DetParam const & theDetParam) const;              // is the det flipped or not?
    
    //---------------------------------------------------------------------------
    //  Cluster-level services.
diff --git a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
index 5e206d9d5bd0b..9b8924988e848 100644
--- a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
+++ b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
@@ -19,7 +19,9 @@ class PixelCPEFast final : public PixelCPEBase
 public:
    struct ClusterParamGeneric : ClusterParam
    {
+      ClusterParamGeneric() {}
       ClusterParamGeneric(const SiPixelCluster & cl) : ClusterParam(cl){}
+
       // The truncation value pix_maximum is an angle-dependent cutoff on the
       // individual pixel signals. It should be applied to all pixels in the
       // cluster [signal_i = fminf(signal_i, pixmax)] before the column and row
@@ -52,6 +54,8 @@ class PixelCPEFast final : public PixelCPEBase
    
    LocalPoint localPosition (DetParam const & theDetParam, ClusterParam & theClusterParam) const override;
    LocalError localError   (DetParam const & theDetParam, ClusterParam & theClusterParam) const override;
+
+   void errorFromTemplates(DetParam const & theDetParam, ClusterParamGeneric & theClusterParam, float qclus) const;
    
    static void
    collect_edge_charges(ClusterParam & theClusterParam,  //!< input, the cluster
diff --git a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h
index 9697470ffb0be..fa326865ced73 100644
--- a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h
+++ b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h
@@ -37,6 +37,8 @@ namespace pixelCPEforGPU {
 
     float x0,y0,z0;  // the vertex in the local coord of the detector
 
+    float sx[3], sy[3]; // the errors...
+
     Frame frame;
   };
 
@@ -208,18 +210,18 @@ namespace pixelCPEforGPU {
   }
 
   constexpr inline
-  void error(CommonParams const & __restrict__ comParams, DetParams const & __restrict__ detParams, ClusParams & cp, uint32_t ic) {
+  void errorFromSize(CommonParams const & __restrict__ comParams, DetParams const & __restrict__ detParams, ClusParams & cp, uint32_t ic) {
     // Edge cluster errors
     cp.xerr[ic]= 0.0050;
     cp.yerr[ic]= 0.0085;
 
     // FIXME these are errors form Run1
     constexpr float xerr_barrel_l1[] = { 0.00115, 0.00120, 0.00088 };
-    constexpr float xerr_barrel_l1_def = 0.01030;
+    constexpr float xerr_barrel_l1_def = 0.00200;           // 0.01030;
     constexpr float yerr_barrel_l1[] = { 0.00375, 0.00230, 0.00250, 0.00250, 0.00230, 0.00230, 0.00210, 0.00210, 0.00240 };
     constexpr float yerr_barrel_l1_def = 0.00210;
     constexpr float xerr_barrel_ln[] = { 0.00115, 0.00120, 0.00088 };
-    constexpr float xerr_barrel_ln_def = 0.01030;
+    constexpr float xerr_barrel_ln_def = 0.00200; // 0.01030;
     constexpr float yerr_barrel_ln[] = { 0.00375, 0.00230, 0.00250, 0.00250, 0.00230, 0.00230, 0.00210, 0.00210, 0.00240 };
     constexpr float yerr_barrel_ln_def = 0.00210;
     constexpr float xerr_endcap[] = { 0.0020, 0.0020 };
@@ -227,12 +229,18 @@ namespace pixelCPEforGPU {
     constexpr float yerr_endcap[] = { 0.00210 };
     constexpr float yerr_endcap_def = 0.00210;
 
+    auto sx = cp.maxRow[ic] - cp.minRow[ic];
+    auto sy = cp.maxCol[ic] - cp.minCol[ic];
+
     // is edgy ?
     bool isEdgeX = cp.minRow[ic] == 0 or cp.maxRow[ic] == phase1PixelTopology::lastRowInModule;
     bool isEdgeY = cp.minCol[ic] == 0 or cp.maxCol[ic] == phase1PixelTopology::lastColInModule;
+    // is one and big?
+    bool isBig1X = (0==sx)  && phase1PixelTopology::isBigPixX(cp.minRow[ic]);
+    bool isBig1Y = (0==sy)  && phase1PixelTopology::isBigPixY(cp.minCol[ic]);
 
-    if (not isEdgeX) {
-      auto sx = cp.maxRow[ic] - cp.minRow[ic];
+
+    if (!isEdgeX && !isBig1X ) {
       if (not detParams.isBarrel) {
         cp.xerr[ic] = sx <std::size(xerr_endcap) ? xerr_endcap[sx] : xerr_endcap_def;
       } else if (detParams.layer == 1) {
@@ -242,8 +250,7 @@ namespace pixelCPEforGPU {
       }
     }
 
-    if (not isEdgeY) {
-      auto sy = cp.maxCol[ic] - cp.minCol[ic];;
+    if (!isEdgeY && !isBig1Y) {
       if (not detParams.isBarrel) {
         cp.yerr[ic] = sy <std::size(yerr_endcap) ? yerr_endcap[sy] : yerr_endcap_def;
       } else if (detParams.layer == 1) {
@@ -254,6 +261,29 @@ namespace pixelCPEforGPU {
     }
   }
 
+
+  constexpr inline
+  void errorFromDB(CommonParams const & __restrict__ comParams, DetParams const & __restrict__ detParams, ClusParams & cp, uint32_t ic) {
+    // Edge cluster errors
+    cp.xerr[ic]= 0.0050f;
+    cp.yerr[ic]= 0.0085f;
+
+    auto sx = cp.maxRow[ic] - cp.minRow[ic];
+    auto sy = cp.maxCol[ic] - cp.minCol[ic];
+
+    // is edgy ?
+    bool isEdgeX = cp.minRow[ic] == 0 or cp.maxRow[ic] == phase1PixelTopology::lastRowInModule;
+    bool isEdgeY = cp.minCol[ic] == 0 or cp.maxCol[ic] == phase1PixelTopology::lastColInModule;
+    // is one and big?
+    uint32_t ix = (0==sx);
+    uint32_t iy = (0==sy);
+    ix+= (0==sx)  && phase1PixelTopology::isBigPixX(cp.minRow[ic]);
+    iy+= (0==sy)  && phase1PixelTopology::isBigPixY(cp.minCol[ic]);
+
+    if (not isEdgeX) cp.xerr[ic] = detParams.sx[ix];
+    if (not isEdgeY) cp.yerr[ic] = detParams.sy[iy];
+  }
+
 }
 
 #endif  // RecoLocalTracker_SiPixelRecHits_pixelCPEforGPU_h
diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
index 4a8c9d945e5f1..6864a046bf1dc 100644
--- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
+++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
@@ -126,7 +126,7 @@ namespace gpuPixelRecHits {
     assert(h < 2000*256);
 
     pixelCPEforGPU::position(cpeParams->commonParams(), cpeParams->detParams(me), clusParams, ic);
-    pixelCPEforGPU::error(cpeParams->commonParams(), cpeParams->detParams(me), clusParams, ic);
+    pixelCPEforGPU::errorFromDB(cpeParams->commonParams(), cpeParams->detParams(me), clusParams, ic);
 
     chargeh[h] = clusParams.charge[ic];
 
@@ -135,8 +135,8 @@ namespace gpuPixelRecHits {
     xl[h]= clusParams.xpos[ic];   
     yl[h]= clusParams.ypos[ic]; 
 
-    xe[h]= clusParams.xerr[ic];
-    ye[h]= clusParams.yerr[ic];
+    xe[h]= clusParams.xerr[ic]*clusParams.xerr[ic];
+    ye[h]= clusParams.yerr[ic]*clusParams.yerr[ic];
     mr[h]= clusParams.minRow[ic];
     mc[h]= clusParams.minCol[ic];
   
diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc
index af7dd7337084e..eb51dd5a2eaeb 100644
--- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc
+++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc
@@ -125,7 +125,84 @@ void PixelCPEFast::fillParamsForGpu() {
     auto vv = p.theDet->surface().position();
     auto rr = pixelCPEforGPU::Rotation(p.theDet->surface().rotation());
     g.frame = pixelCPEforGPU::Frame(vv.x(),vv.y(),vv.z(),rr);
-  }
+
+
+    // errors .....
+   ClusterParamGeneric cp;
+   auto gvx = p.theOrigin.x() + 40.f*m_commonParamsGPU.thePitchX;
+   auto gvy = p.theOrigin.y();
+   auto gvz = 1.f/p.theOrigin.z();
+   //--- Note that the normalization is not required as only the ratio used
+
+   // calculate angles
+   cp.cotalpha = gvx*gvz;
+   cp.cotbeta  = gvy*gvz;
+
+   cp.with_track_angle = false;
+
+   auto lape = p.theDet->localAlignmentError();
+   if ( lape.invalid() ) lape = LocalError(); // zero....
+
+#ifdef DUMP_ERRORS   
+   auto m=10000.f;
+   for (float qclus = 15000; qclus<35000; qclus+=15000){
+     errorFromTemplates(p,cp,qclus);
+
+     std::cout << i << ' ' << qclus << ' ' << cp.pixmx
+               << ' ' << m*cp.sigmax << ' ' << m*cp.sx1 << ' ' << m*cp.sx2 
+               << ' ' << m*cp.sigmay << ' ' << m*cp.sy1 << ' ' << m*cp.sy2
+              << std::endl;
+   }
+   std::cout << i << ' ' << m*std::sqrt(lape.xx())  <<' '<< m*std::sqrt(lape.yy()) << std::endl;
+#endif   
+
+   
+   errorFromTemplates(p,cp,20000.f);
+   g.sx[0] = cp.sigmax;
+   g.sx[1] = cp.sx1;
+   g.sx[2] = cp.sx2;
+
+   g.sy[0] = cp.sigmay;
+   g.sy[1] = cp.sy1;
+   g.sy[2] = cp.sy2;
+   
+   
+   /*
+    // from run1??
+    if (i<96) {
+      g.sx[0] = 0.00120;
+      g.sx[1] = 0.00115;
+      g.sx[2] = 0.0050;
+
+      g.sy[0] = 0.00210;
+      g.sy[1] = 0.00375;
+      g.sy[2] = 0.0085;
+    } else if (g.isBarrel) {
+      g.sx[0] = 0.00120;
+      g.sx[1] = 0.00115;
+      g.sx[2] = 0.0050;
+
+      g.sy[0] = 0.00210;
+      g.sy[1] = 0.00375;
+      g.sy[2] = 0.0085;
+   } else {
+      g.sx[0] = 0.0020;
+      g.sx[1] = 0.0020;
+      g.sx[2] = 0.0050;
+
+      g.sy[0] = 0.0021;
+      g.sy[1] = 0.0021;
+      g.sy[2] = 0.0085;
+   }
+   */
+   
+
+   for (int i=0; i<3; ++i) {
+     g.sx[i] = std::sqrt(g.sx[i]*g.sx[i]+lape.xx());
+     g.sy[i] = std::sqrt(g.sy[i]*g.sy[i]+lape.yy());
+   }
+
+ }
 }
 
 PixelCPEFast::~PixelCPEFast() {}
@@ -143,25 +220,15 @@ PixelCPEBase::ClusterParam* PixelCPEFast::createClusterParam(const SiPixelCluste
    return new ClusterParamGeneric(cl);
 }
 
-//-----------------------------------------------------------------------------
-//! Hit position in the local frame (in cm).  Unlike other CPE's, this
-//! one converts everything from the measurement frame (in channel numbers)
-//! into the local frame (in centimeters).
-//-----------------------------------------------------------------------------
-LocalPoint
-PixelCPEFast::localPosition(DetParam const & theDetParam, ClusterParam & theClusterParamBase) const
-{
-   ClusterParamGeneric & theClusterParam = static_cast<ClusterParamGeneric &>(theClusterParamBase);
 
-   assert(!theClusterParam.with_track_angle); 
-   
-   if ( UseErrorsFromTemplates_ ) {
-      
-      float qclus = theClusterParam.theCluster->charge();
+
+void
+PixelCPEFast::errorFromTemplates(DetParam const & theDetParam, ClusterParamGeneric & theClusterParam, float qclus) const
+{
       float locBz = theDetParam.bz;
       float locBx = theDetParam.bx;
       //cout << "PixelCPEFast::localPosition(...) : locBz = " << locBz << endl;
-      
+
       theClusterParam.pixmx  = std::numeric_limits<int>::max();  // max pixel charge for truncation of 2-D cluster
 
       theClusterParam.sigmay = -999.9; // CPE Generic y-error for multi-pixel cluster
@@ -170,28 +237,43 @@ PixelCPEFast::localPosition(DetParam const & theDetParam, ClusterParam & theClus
       theClusterParam.sy2    = -999.9; // CPE Generic y-error for single double-pixel cluster
       theClusterParam.sx1    = -999.9; // CPE Generic x-error for single single-pixel cluster
       theClusterParam.sx2    = -999.9; // CPE Generic x-error for single double-pixel cluster
-      
+
       float dummy;
-      
+
       SiPixelGenError gtempl(thePixelGenError_);
       int gtemplID_ = theDetParam.detTemplateId;
-      
-      theClusterParam.qBin_ = gtempl.qbin( gtemplID_, theClusterParam.cotalpha, theClusterParam.cotbeta, locBz, locBx, qclus, 
+
+      theClusterParam.qBin_ = gtempl.qbin( gtemplID_, theClusterParam.cotalpha, theClusterParam.cotbeta, locBz, locBx, qclus,
                                           false,
                                           theClusterParam.pixmx, theClusterParam.sigmay, dummy,
                                           theClusterParam.sigmax, dummy, theClusterParam.sy1,
                                           dummy, theClusterParam.sy2, dummy, theClusterParam.sx1,
                                           dummy, theClusterParam.sx2, dummy );
-      
+
       theClusterParam.sigmax = theClusterParam.sigmax * micronsToCm;
       theClusterParam.sx1 = theClusterParam.sx1 * micronsToCm;
       theClusterParam.sx2 = theClusterParam.sx2 * micronsToCm;
-      
+
       theClusterParam.sigmay = theClusterParam.sigmay * micronsToCm;
       theClusterParam.sy1 = theClusterParam.sy1 * micronsToCm;
       theClusterParam.sy2 = theClusterParam.sy2 * micronsToCm;
-      
-   } // if ( UseErrorsFromTemplates_ )
+}
+
+//-----------------------------------------------------------------------------
+//! Hit position in the local frame (in cm).  Unlike other CPE's, this
+//! one converts everything from the measurement frame (in channel numbers)
+//! into the local frame (in centimeters).
+//-----------------------------------------------------------------------------
+LocalPoint
+PixelCPEFast::localPosition(DetParam const & theDetParam, ClusterParam & theClusterParamBase) const
+{
+   ClusterParamGeneric & theClusterParam = static_cast<ClusterParamGeneric &>(theClusterParamBase);
+
+   assert(!theClusterParam.with_track_angle); 
+   
+   if ( UseErrorsFromTemplates_ ) {
+      errorFromTemplates(theDetParam, theClusterParam, theClusterParam.theCluster->charge());
+   } 
    else {
      theClusterParam.qBin_ = 0;
    }