From 7501972a56ac28835c6762fd653ae86f1193dd4d Mon Sep 17 00:00:00 2001 From: Vincenzo Innocente Date: Fri, 29 Jun 2018 10:21:35 +0200 Subject: [PATCH] Migrated PixelRecHit to Heterogeneous producer (cms-patatrack#81) Migrate PixelRecHit EDProducer to HeterogeneousEDProducer, including the cpu product. Data structures on gpu now include everything needed for Doublets, CA and fit. Layer splitting done: phi sorting (or partial sorting) requires #69. Includes some cleanup and bug fixes. --- .../interface/phase1PixelTopology.h | 9 ++++ .../python/RecoLocalTracker_cff.py | 9 +++- .../SiPixelRecHits/plugins/gpuPixelRecHits.h | 45 +++++++++++++------ .../python/SiPixelRecHits_cfi.py | 3 -- 4 files changed, 48 insertions(+), 18 deletions(-) diff --git a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h index 455de58ce3408..37c97a92a3eaa 100644 --- a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h +++ b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h @@ -20,6 +20,15 @@ namespace phase1PixelTopology { constexpr uint32_t numPixsInModule = uint32_t(numRowsInModule)* uint32_t(numColsInModule); + constexpr uint32_t numberOfModules = 1856; + + constexpr uint32_t layerStart[11] = {0,96,320,672,1184,1296,1408,1520,1632,1744,1856}; + constexpr char const * layerName[10] = {"BL1","BL2","BL3","BL4", + "E+1", "E+2", "E+3", + "E-1", "E-2", "E-3" + }; + + // this is for the ROC n<512 (upgrade 1024) constexpr inline uint16_t divu52(uint16_t n) { diff --git a/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py b/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py index b601e310db645..6d803d40bc870 100644 --- a/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py +++ b/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py @@ -17,16 +17,23 @@ striptrackerlocalreco = cms.Sequence(siStripZeroSuppression*siStripClusters*siStripMatchedRecHits) trackerlocalreco = cms.Sequence(pixeltrackerlocalreco*striptrackerlocalreco*clusterSummaryProducer) + from RecoLocalTracker.SiPixelClusterizer.siPixelClustersHeterogeneous_cfi import * from RecoLocalTracker.SiPixelClusterizer.siPixelFedCablingMapGPUWrapper_cfi import * from CalibTracker.SiPixelESProducers.siPixelGainCalibrationForHLTGPU_cfi import * +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import * +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneousConverter_cfi import siPixelRecHitHeterogeneousConverter as _siPixelRecHitHeterogeneousConverter +gpu.toReplaceWith(siPixelRecHitsPreSplitting, _siPixelRecHitHeterogeneousConverter.clone()) + + + from Configuration.ProcessModifiers.gpu_cff import gpu _pixeltrackerlocalreco_gpu = pixeltrackerlocalreco.copy() _pixeltrackerlocalreco_gpu.replace(siPixelClustersPreSplitting, siPixelClustersHeterogeneous+siPixelClustersPreSplitting) +_pixeltrackerlocalreco_gpu.replace(siPixelRecHitsPreSplitting, siPixelRecHitHeterogeneous+siPixelRecHitsPreSplitting) gpu.toReplaceWith(pixeltrackerlocalreco, _pixeltrackerlocalreco_gpu) - from RecoLocalTracker.SiPhase2Clusterizer.phase2TrackerClusterizer_cfi import * from RecoLocalTracker.Phase2TrackerRecHits.Phase2StripCPEGeometricESProducer_cfi import * diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index 98a5198232591..c0e7841658e93 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -6,10 +6,14 @@ #include #include +#include "DataFormats/Math/interface/approx_atan2.h" #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h" namespace gpuPixelRecHits { + + + // to be moved in common namespace... constexpr uint16_t InvId=9999; // must be > MaxNumModules @@ -20,6 +24,7 @@ namespace gpuPixelRecHits { __global__ void getHits(pixelCPEforGPU::ParamsOnGPU const * cpeParams, + float const * bs, uint16_t const * id, uint16_t const * x, uint16_t const * y, @@ -27,13 +32,15 @@ namespace gpuPixelRecHits { uint32_t const * digiModuleStart, uint32_t const * clusInModule, uint32_t const * moduleId, - int32_t const * clus, + int32_t const * clus, int numElements, uint32_t const * hitsModuleStart, int32_t * chargeh, - float * xh, float * yh, float * zh, - float * xe, float * ye, uint16_t * mr, - bool local) // if true fill just x & y in local coord + uint16_t * detInd, + float * xg, float * yg, float * zg, float * rg, int16_t * iph, + float * xl, float * yl, + float * xe, float * ye, + uint16_t * mr, uint16_t * mc) { // as usual one block per module __shared__ ClusParams clusParams; @@ -108,16 +115,26 @@ namespace gpuPixelRecHits { chargeh[h] = clusParams.charge[ic]; - if (local) { - xh[h] = clusParams.xpos[ic]; - yh[h] = clusParams.ypos[ic]; - } else { - cpeParams->detParams(me).frame.toGlobal(clusParams.xpos[ic], clusParams.ypos[ic], - xh[h], yh[h], zh[h] ); - } - xe[h] = clusParams.xerr[ic]; - ye[h] = clusParams.yerr[ic]; - mr[h] = clusParams.minRow[ic]; + detInd[h] = me; + + xl[h]= clusParams.xpos[ic]; + yl[h]= clusParams.ypos[ic]; + + xe[h]= clusParams.xerr[ic]; + ye[h]= clusParams.yerr[ic]; + mr[h]= clusParams.minRow[ic]; + mc[h]= clusParams.minCol[ic]; + + // to global and compute phi... + cpeParams->detParams(me).frame.toGlobal(xl[h],yl[h], xg[h],yg[h],zg[h]); + // here correct for the beamspot... + xg[h]-=bs[0]; + yg[h]-=bs[1]; + zg[h]-=bs[2]; + + rg[h] = std::sqrt(xg[h]*xg[h]+yg[h]*yg[h]); + iph[h] = unsafe_atan2s<7>(yg[h],xg[h]); + } } diff --git a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py index 87a0d3118554b..5844235e29596 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py +++ b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py @@ -7,9 +7,6 @@ VerboseLevel = cms.untracked.int32(0), ) -from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import siPixelRecHitHeterogeneous as _siPixelRecHitHeterogeneous -gpu.toReplaceWith(siPixelRecHits, _siPixelRecHitHeterogeneous) - siPixelRecHitsPreSplitting = siPixelRecHits.clone( src = 'siPixelClustersPreSplitting' )