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/Configuration/python/customizePixelTracksForProfiling.py

@@ -1,6 +1,41 @@
 import FWCore.ParameterSet.Config as cms
 
-def customizePixelTracksForProfiling(process):
+def customizePixelTracksForProfilingGPUOnly(process):
+    process.MessageLogger.cerr.FwkReport.reportEvery = 100
+
+    process.Raw2Hit = cms.Path(process.offlineBeamSpot+process.offlineBeamSpotCUDA+process.siPixelClustersCUDAPreSplitting+process.siPixelRecHitsCUDAPreSplitting)
+
+    process.load('RecoPixelVertexing.PixelTriplets.caHitNtupletCUDA_cfi')
+    process.load('RecoPixelVertexing.PixelVertexFinding.pixelVertexCUDA_cfi')
+    process.TVreco = cms.Path(process.caHitNtupletCUDA+process.pixelVertexCUDA)
+
+    process.schedule = cms.Schedule(process.Raw2Hit, process.TVreco)
+    return process
+
+def customizePixelTracksForProfilingSoAonCPU(process):
+    process = customizePixelTracksForProfilingGPUOnly(process)
+
+    process.pixelVertexSoA = process.pixelVertexCUDA.clone()
+    process.pixelVertexSoA.onGPU = False
+    process.pixelVertexSoA.pixelTrackSrc = 'pixelTrackSoA'
+    process.TVSoAreco = cms.Path(process.caHitNtupletCUDA+process.pixelTrackSoA+process.pixelVertexSoA)
+
+    process.schedule = cms.Schedule(process.Raw2Hit, process.TVSoAreco)
+
+    return process
+
+def customizePixelTracksForProfilingEnableTransfer(process):
+    process = customizePixelTracksForProfilingGPUOnly(process)
+
+    process.load('RecoPixelVertexing.PixelTrackFitting.pixelTrackSoA_cfi')
+    process.load('RecoPixelVertexing.PixelVertexFinding.pixelVertexSoA_cfi')
+    process.toSoA = cms.Path(process.pixelTrackSoA+process.pixelVertexSoA)
+
+    process.schedule = cms.Schedule(process.Raw2Hit, process.TVreco, process.toSoA)
+    return process
+
+def customizePixelTracksForProfilingEnableConversion(process):
+    # use old trick of output path
     process.MessageLogger.cerr.FwkReport.reportEvery = 100
 
     process.out = cms.OutputModule("AsciiOutputModule",
@@ -17,21 +52,3 @@ def customizePixelTracksForProfiling(process):
 
     return process
 
-def customizePixelTracksForProfilingDisableConversion(process):
-    process = customizePixelTracksForProfiling(process)
-
-    # Disable conversions to legacy
-    process.pixelTracksHitQuadruplets.gpuEnableConversion = False
-    process.pixelTracks.gpuEnableConversion = False
-    process.pixelVertices.gpuEnableConversion = False
-
-    return process
-
-def customizePixelTracksForProfilingDisableTransfer(process):
-    process = customizePixelTracksForProfilingDisableConversion(process)
-
-    # Disable "unnecessary" transfers to CPU
-    process.pixelTracksHitQuadruplets.gpuEnableTransfer = False
-    process.pixelVertices.gpuEnableTransfer = False
-
-    return process

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)