Port the whole pixel workflow to new heterogeneous framework (#384)

  - port the whole pixel workflow to new heterogeneous framework
  - implement a legacy cluster to SoA converter for the pixel RecHits
  - update the vertex producer to run on CPU as well as GPU

CUDADataFormats/Vertex/BuildFile.xml

@@ -0,0 +1,10 @@
+<use name="cuda-api-wrappers"/>
+<use name="rootcore"/>
+<use name="FWCore/ServiceRegistry"/>
+<use name="FWCore/ParameterSetReader"/>
+<use name="HeterogeneousCore/CUDAServices"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+<use name="eigen"/>
+<export>
+    <lib name="1"/>
+</export>

CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h

@@ -0,0 +1,15 @@
+#ifndef CUDADataFormatsVertexZVertexHeterogeneous_H
+#define CUDADataFormatsVertexZVertexHeterogeneous_H
+
+#include "CUDADataFormats/Vertex/interface/ZVertexSoA.h"
+#include "CUDADataFormats/Common/interface/HeterogeneousSoA.h"
+#include "CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h"
+
+
+using ZVertexHeterogeneous = HeterogeneousSoA<ZVertexSoA>;
+#ifndef __CUDACC__
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+using ZVertexCUDAProduct =  CUDAProduct<ZVertexHeterogeneous>;
+#endif
+
+#endif

CUDADataFormats/Vertex/interface/ZVertexSoA.h

@@ -0,0 +1,29 @@
+#ifndef CUDADataFormatsVertexZVertexSoA_H
+#define CUDADataFormatsVertexZVertexSoA_H
+
+#include<cstdint>
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
+
+
+// SOA for vertices
+// These vertices are clusterized and fitted only along the beam line (z)
+// to obtain their global coordinate the beam spot position shall be added (eventually correcting for the beam angle as well)
+struct ZVertexSoA {
+    static constexpr uint32_t MAXTRACKS = 32*1024;
+    static constexpr uint32_t MAXVTX = 1024;
+
+    int16_t idv[MAXTRACKS];    // vertex index for each associated (original) track  (-1 == not associate)
+    float zv[MAXVTX];          // output z-posistion of found vertices
+    float wv[MAXVTX];          // output weight (1/error^2) on the above
+    float chi2[MAXVTX];        // vertices chi2
+    float ptv2[MAXVTX];        // vertices pt^2
+    int32_t ndof[MAXVTX];      // vertices number of dof (reused as workspace for the number of nearest neighbours)
+    uint16_t sortInd[MAXVTX];  // sorted index (by pt2)  ascending
+    uint32_t nvFinal;          // the number of vertices
+
+    __host__ __device__ void init() { nvFinal = 0; }
+
+};
+
+#endif // CUDADataFormatsVertexZVertexSoA.H
+

CUDADataFormats/Vertex/src/classes.h

@@ -0,0 +1,8 @@
+#ifndef CUDADataFormats__src_classes_h
+#define CUDADataFormats__src_classes_h
+
+#include "CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h"
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif  

CUDADataFormats/Vertex/src/classes_def.xml

@@ -0,0 +1,6 @@
+<lcgdict>
+  <class name="CUDAProduct<ZVertexHeterogeneous>" persistent="false"/>
+  <class name="edm::Wrapper<ZVertexCUDAProduct>" persistent="false"/>
+  <class name="ZVertexHeterogeneous" persistent="false"/>
+  <class name="edm::Wrapper<ZVertexHeterogeneous>" persistent="false"/>
+</lcgdict>

RecoPixelVertexing/Configuration/python/RecoPixelVertexing_cff.py

@@ -8,3 +8,18 @@
 #from RecoVertex.PrimaryVertexProducer.OfflinePixel3DPrimaryVertices_cfi import *
 recopixelvertexingTask = cms.Task(pixelTracksTask,pixelVertices)
 recopixelvertexing = cms.Sequence(recopixelvertexingTask)
+
+from Configuration.ProcessModifiers.gpu_cff import gpu
+
+from RecoPixelVertexing.PixelVertexFinding.pixelVertexCUDA_cfi import pixelVertexCUDA
+from RecoPixelVertexing.PixelVertexFinding.pixelVertexSoA_cfi import pixelVertexSoA
+from RecoPixelVertexing.PixelVertexFinding.pixelVertexFromSoA_cfi import pixelVertexFromSoA as _pixelVertexFromSoA
+
+_pixelVertexingCUDATask = cms.Task(pixelTracksTask,pixelVertexCUDA,pixelVertexSoA,pixelVertices)
+
+# pixelVertexSoAonCPU = pixelVertexCUDA.clone()
+# pixelVertexSoAonCPU.onGPU = False;
+
+gpu.toReplaceWith(pixelVertices,_pixelVertexFromSoA)
+gpu.toReplaceWith(recopixelvertexingTask,_pixelVertexingCUDATask)
+

RecoPixelVertexing/PixelVertexFinding/python/PixelVertexes_cfi.py

@@ -18,8 +18,3 @@
        refToPSet_ = cms.string('pvClusterComparer')
     )
 )
-
-
-from Configuration.ProcessModifiers.gpu_cff import gpu
-from RecoPixelVertexing.PixelVertexFinding.pixelVertexHeterogeneousProducer_cfi import pixelVertexHeterogeneousProducer as _pixelVertexHeterogeneousProducer
-gpu.toReplaceWith(pixelVertices, _pixelVertexHeterogeneousProducer)