digital calibrator, concurrent clusterizer (optimized), CPE and RecHIts

CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationServiceBase.h

@@ -73,6 +73,8 @@ class SiPixelGainCalibrationServicePayloadGetter : public SiPixelGainCalibration
 
   void    setESObjects(const edm::EventSetup& es ) override;
 
+  thePayloadObject const & payload() const { return *ped; }
+
   std::vector<uint32_t> getDetIds() override;
   double getGainLow() override;
   double getGainHigh() override;

CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h

@@ -73,6 +73,9 @@ class SiPixelGainCalibrationForHLT {
   float getPedLow() const { return minPed_; }
   float getPedHigh() const { return maxPed_; }
 
+  std::vector<char> const & data() const { return v_pedestals;}
+  std::vector<DetRegistry> const & getIndexes() const { return indexes; }
+
   // Set and get public methods
   void  setData(float ped, float gain, std::vector<char>& vped, bool thisColumnIsDead = false, bool thisColumnIsNoisy = false);
   void  setDeadColumn(const int& nRows, std::vector<char>& vped)  { setData(0, 0 /*dummy values, not used*/, vped, true, false); }
@@ -84,7 +87,7 @@ class SiPixelGainCalibrationForHLT {
   float getPed   (const int& col, const int& row, const Range& range, const int& nCols, bool& isDeadColumn, bool& isNoisyColumn ) const;
   float getGain  (const int& col, const int& row, const Range& range, const int& nCols, bool& isDeadColumn, bool& isNoisyColumn ) const;
 
-  private:
+private:
 
   float   encodeGain(const float& gain);
   float   encodePed (const float& ped);

CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h

@@ -0,0 +1,69 @@
+#pragma once
+
+#include<cstdint>
+#include<tuple>
+#include<cassert>
+#include<cstdio>
+
+struct SiPixelGainForHLTonGPU_DecodingStructure{
+  uint8_t gain;
+  uint8_t ped;
+};
+
+
+// copy of SiPixelGainCalibrationForHLT
+class SiPixelGainForHLTonGPU {
+
+ public:
+
+  using DecodingStructure = SiPixelGainForHLTonGPU_DecodingStructure;
+
+  using Range = std::pair<uint32_t,uint32_t>;
+
+
+  inline __host__ __device__
+  std::pair<float,float> getPedAndGain(uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn ) const {
+
+
+    auto range = rangeAndCols[moduleInd].first;
+    auto nCols = rangeAndCols[moduleInd].second;
+
+    // determine what averaged data block we are in (there should be 1 or 2 of these depending on if plaquette is 1 by X or 2 by X
+    unsigned int lengthOfColumnData  = (range.second-range.first)/nCols;
+    unsigned int lengthOfAveragedDataInEachColumn = 2;  // we always only have two values per column averaged block 
+    unsigned int numberOfDataBlocksToSkip = row / numberOfRowsAveragedOver_;
+
+
+    auto offset = range.first + col*lengthOfColumnData + lengthOfAveragedDataInEachColumn*numberOfDataBlocksToSkip;
+
+    assert(offset<range.second);
+    assert(offset<3088384);
+    assert(0==offset%2);
+
+    auto s = v_pedestals[offset/2];
+
+    isDeadColumn = (s.ped & 0xFF) == deadFlag_;
+    isNoisyColumn = (s.ped & 0xFF) == noisyFlag_;
+
+    return std::make_pair(decodePed(s.ped & 0xFF),decodeGain(s.gain & 0xFF));
+
+  }
+
+
+
+  constexpr float decodeGain(unsigned int gain) const {return gain*gainPrecision + minGain_;}
+  constexpr float decodePed (unsigned int ped) const { return ped*pedPrecision + minPed_;}
+
+  DecodingStructure * v_pedestals;
+  std::pair<Range, int> rangeAndCols[2000];
+
+  float  minPed_, maxPed_, minGain_, maxGain_;
+
+  float pedPrecision, gainPrecision;
+
+  unsigned int numberOfRowsAveragedOver_; // this is 80!!!!
+  unsigned int nBinsToUseForEncoding_;
+  unsigned int deadFlag_;
+  unsigned int noisyFlag_;
+};
+

DataFormats/GeometrySurface/interface/SOARotation.h

@@ -0,0 +1,160 @@
+#pragma once
+
+template <class T> class TkRotation;
+
+// to be moved in an external common library???
+
+/** Rotation matrix used by SOA (as in GPU 
+ */
+
+template <class T>
+class SOARotation {
+public:
+
+  constexpr inline
+  SOARotation(){}
+
+  constexpr inline
+  explicit SOARotation(T) : 
+    R11( 1), R12( 0), R13( 0), 
+    R21( 0), R22( 1), R23( 0),
+    R31( 0), R32( 0), R33( 1) {}
+
+  constexpr inline
+  SOARotation( T xx, T xy, T xz, T yx, T yy, T yz, T zx, T zy, T zz) :
+    R11(xx), R12(xy), R13(xz), 
+    R21(yx), R22(yy), R23(yz),
+    R31(zx), R32(zy), R33(zz) {}
+
+  constexpr inline
+  SOARotation( const T* p) :
+    R11(p[0]), R12(p[1]), R13(p[2]), 
+    R21(p[3]), R22(p[4]), R23(p[5]),
+    R31(p[6]), R32(p[7]), R33(p[8]) {}
+
+
+  template <typename U>
+  constexpr inline
+  SOARotation( const TkRotation<U>& a) : 
+    R11(a.xx()), R12(a.xy()), R13(a.xz()), 
+    R21(a.yx()), R22(a.yy()), R23(a.yz()),
+    R31(a.zx()), R32(a.zy()), R33(a.zz()) {}
+
+  constexpr inline
+  SOARotation transposed() const {
+      return SOARotation( R11, R21, R31, 
+			 R12, R22, R32,
+			 R13, R23, R33);
+  }
+
+  // if frame this is to local
+   constexpr inline
+ void multiply(T const vx, T const vy, T const vz,
+              T & ux, T & uy, T & uz
+             ) const {
+     ux = R11*vx + R12*vy + R13*vz;
+     uy = R21*vx + R22*vy + R23*vz;
+     uz = R31*vx + R32*vy + R33*vz;
+  }
+
+ // if frame this is to global
+ constexpr inline
+ void multiplyInverse (T const vx, T const vy, T const vz,
+              T & ux, T & uy, T & uz
+             ) const {
+     ux = R11*vx + R21*vy + R31*vz;
+     uy = R12*vx + R22*vy + R32*vz;  
+     uz = R13*vx + R23*vy + R33*vz;
+  }
+
+
+ // if frame this is to global
+ constexpr inline
+ void multiplyInverse (T const vx, T const vy,
+              T & ux, T & uy, T & uz
+             ) const {
+     ux = R11*vx + R21*vy;
+     uy = R12*vx + R22*vy;
+     uz = R13*vx + R23*vy;
+  }
+
+
+  constexpr inline
+  T const &xx() const { return R11;} 
+  constexpr inline
+  T const &xy() const { return R12;} 
+  constexpr inline
+  T const &xz() const { return R13;} 
+  constexpr inline
+  T const &yx() const { return R21;} 
+  constexpr inline
+  T const &yy() const { return R22;} 
+  constexpr inline
+  T const &yz() const { return R23;} 
+  constexpr inline
+  T const &zx() const { return R31;} 
+  constexpr inline
+  T const &zy() const { return R32;} 
+  constexpr inline
+  T const &zz() const { return R33;} 
+
+private:
+
+  T R11, R12, R13;
+  T R21, R22, R23;
+  T R31, R32, R33;
+};
+
+
+template <class T>
+class SOAFrame {
+public:
+
+  constexpr inline   
+  SOAFrame(){}
+
+  constexpr inline
+  SOAFrame(T ix, T iy, T iz, SOARotation<T> const & irot) :
+     px(ix),py(iy),pz(iz), rot(irot){}
+
+  constexpr inline
+  SOARotation<T> const & rotation() const { return rot;}
+
+  constexpr inline
+  void toLocal(T const vx, T const vy,	T const vz,
+               T & ux, T & uy, T & uz
+              ) const {
+      rot.multiply(vx-px,vy-py,vz-pz,ux,uy,uz);
+  }
+
+
+  constexpr inline
+  void toGlobal(T const vx, T const vy, T const vz,
+                T & ux, T & uy, T & uz
+               ) const {
+       rot.multiplyInverse(vx,vy,vz,ux,uy,uz);
+       ux+=px; uy+=py; uz+=pz; 
+  }
+
+  constexpr inline
+  void toGlobal(T const vx, T const vy,
+                T & ux, T & uy, T & uz
+             ) const {
+       rot.multiplyInverse(vx,vy,ux,uy,uz);
+       ux+=px; uy+=py; uz+=pz;
+  }
+
+
+  constexpr inline 
+  T x() const { return px;} 
+  constexpr inline
+  T y()	const { return py;}
+  constexpr inline
+  T z()	const { return pz;}
+
+private:
+
+  T px, py, pz;
+  SOARotation<T> rot;
+
+};

DataFormats/GeometrySurface/test/BuildFile.xml

@@ -14,3 +14,18 @@
 </bin>
 <bin   file="Bounds_t.cpp">
 </bin>
+
+
+<use name="cuda"/>
+<use name="cuda-api-wrappers"/>
+<bin file="gpuFrameTransformKernel.cu, gpuFrameTransformTest.cpp" name="gpuFrameTransformTest">
+  <flags CXXFLAGS="-g"/>
+  <flags CUDA_FLAGS="-O3 --expt-relaxed-constexpr"/>
+</bin>
+
+
+<bin file="gpuFrameTransformKernel.cu, gpuFrameTransformTest.cpp" name="gpuFrameTransformTestRep">
+  <flags CUDA_FLAGS="-O3 --expt-relaxed-constexpr -fmad=false -ftz=false -prec-div=true -prec-sqrt=true"/>
+  <flags CXXFLAGS="-ffp-contract=off"/>
+</bin>
+

DataFormats/GeometrySurface/test/FrameTransformTest.cpp

@@ -1,4 +1,4 @@
-#include <iostream>
+#include "DataFormats/GeometrySurface/interface/SOARotation.h"
 #include "DataFormats/GeometrySurface/interface/TkRotation.h"
 #include "DataFormats/GeometrySurface/interface/GloballyPositioned.h"
 #include "DataFormats/GeometrySurface/interface/BoundPlane.h"
@@ -7,12 +7,16 @@
 
 #include <cmath>
 
+#include <iostream>
+
 using namespace std;
 template<typename T>
 void  go() {
 
     typedef TkRotation<T>                   Rotation;
+    typedef SOARotation<T>                  SRotation;
     typedef GloballyPositioned<T>           Frame;
+    typedef SOAFrame<T>                     SFrame;
     typedef typename Frame::PositionType             Position;
     typedef typename Frame::GlobalVector             GlobalVector;
     typedef typename Frame::GlobalPoint              GlobalPoint;
@@ -23,6 +27,8 @@ void  go() {
     std::cout << "size of Pos     " << sizeof(Position) << std::endl;
     std::cout << "size of Point   " << sizeof(GlobalPoint) << std::endl;
     std::cout << "size of Frame   " << sizeof(Frame) << std::endl;
+    std::cout << "size of soa Rot     " << sizeof(SRotation) << std::endl;
+    std::cout << "size of soa Frame     " << sizeof(SFrame) << std::endl;
 
     double a = 0.01;
     double ca = cos(a);
@@ -51,15 +57,28 @@ void  go() {
     // Rotation r3 = r2*r1;
     Rotation r3 = r2*r1.transposed();
 
+    SRotation sr3(r3);
 
     GlobalPoint pos2(f2.position());
     LocalPoint lp3 = f1.toLocal(pos2);
     Frame f3( GlobalPoint(lp3.basicVector()), r3);
     cout << "f3.position() " << f3.position() << endl;
     cout << "f3.rotation() " << endl << f3.rotation() << endl;
 
+    SFrame sf1(f1.position().x(),
+               f1.position().y(),
+               f1.position().z(),
+               f1.rotation()
+              );
+
+    SFrame sf3(f3.position().x(),
+               f3.position().y(),
+               f3.position().z(),
+     	       	 sr3
+  	      );
+
 // test
-    GlobalPoint gp( 11,22,33);
+    GlobalPoint gp(11,22,33);
     LocalPoint p_in1 = f1.toLocal( gp);
     typename Frame::ToLocal ff1(f1);
     LocalPoint p_in2 = f2.toLocal( gp);
@@ -69,8 +88,19 @@ void  go() {
     cout << "p_in2 " << p_in2 << endl;
     cout << "p_in3 " << p_in3 << endl;
 
-    LocalPoint p_in1_from3( f3.toGlobal(p_in3).basicVector());
-    cout << "p_in1_from3 + " << p_in1_from3 << endl;
+    LocalPoint p_in1_from3(f3.toGlobal(p_in3).basicVector());
+    cout << "p_in1_from3 " << p_in1_from3 << endl;
+
+    T xx,yy,zz;
+    sf1.toLocal(gp.x(),gp.y(),gp.z(),
+                xx,yy,zz);
+    cout << "p_in1 soa " << xx<<','<<yy<<','<< zz << endl;
+    sf3.toLocal(p_in1.x(),p_in1.y(),p_in1.z(),
+       	       	 xx,yy,zz);
+    cout << "p_in3 soa " << xx<<','<<yy<<','<< zz << endl;
+    sf3.toGlobal(p_in3.x(),p_in3.y(),p_in3.z(),
+                 xx,yy,zz);
+    cout << "p_in1_from3 soa " << xx<<','<<yy<<','<< zz << endl;
 
     BoundPlane plane(f2.position(), f2.rotation());