From cbe13d07fd8bb75ed23d4f72d5f81e4f11dacb69 Mon Sep 17 00:00:00 2001 From: Vincenzo Innocente Date: Wed, 26 Sep 2018 17:53:25 +0200 Subject: [PATCH] Introduce Cluster Charge Cut, optimize Histogram (bucket sorting) (cms-patatrack#171) --- CalibTracker/SiPixelESProducers/BuildFile.xml | 2 + .../SiPixelESProducers/plugins/BuildFile.xml | 2 + .../StandardSequences/python/RawToDigi_cff.py | 3 +- .../beampixel_dqm_sourceclient-live_cfg.py | 230 ++++++++++-------- .../l1tstage2_dqm_sourceclient-live_cfg.py | 11 +- ...tage2emulator_dqm_sourceclient-live_cfg.py | 9 +- .../python/SiPixelRawToDigi_cfi.py | 8 + .../interface/phase1PixelTopology.h | 9 + .../python/RecoLocalTracker_cff.py | 4 + .../SiPixelClusterizer/plugins/BuildFile.xml | 11 +- .../plugins/SiPixelRawToClusterGPUKernel.cu | 23 +- .../plugins/gpuClusterChargeCut.h | 97 ++++++++ .../plugins/gpuClustering.h | 151 +++++++----- .../plugins/gpuClusteringConstants.h | 2 + .../SiPixelClusterizerPreSplitting_cfi.py | 8 +- .../SiPixelClusterizer/test/BuildFile.xml | 15 ++ RecoLocalTracker/SiPixelRecHits/BuildFile.xml | 5 + .../SiPixelRecHits/plugins/BuildFile.xml | 17 +- .../SiPixelRecHits/plugins/PixelRecHits.cu | 14 +- .../SiPixelRecHits/plugins/gpuPixelRecHits.h | 41 +++- .../python/PixelCPEESProducers_cff.py | 1 + 21 files changed, 456 insertions(+), 207 deletions(-) create mode 100644 RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h diff --git a/CalibTracker/SiPixelESProducers/BuildFile.xml b/CalibTracker/SiPixelESProducers/BuildFile.xml index e9d22b32f0afb..69d258da21ed1 100644 --- a/CalibTracker/SiPixelESProducers/BuildFile.xml +++ b/CalibTracker/SiPixelESProducers/BuildFile.xml @@ -7,7 +7,9 @@ + + diff --git a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml b/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml index 44db9d9ba0582..b33657e273036 100644 --- a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml +++ b/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml @@ -6,6 +6,8 @@ + + diff --git a/Configuration/StandardSequences/python/RawToDigi_cff.py b/Configuration/StandardSequences/python/RawToDigi_cff.py index e36243fc32f16..4380d33f134cc 100644 --- a/Configuration/StandardSequences/python/RawToDigi_cff.py +++ b/Configuration/StandardSequences/python/RawToDigi_cff.py @@ -76,7 +76,8 @@ RawToDigi_pixelOnly = cms.Sequence(siPixelDigis) scalersRawToDigi.scalersInputTag = 'rawDataCollector' -siPixelDigis.InputLabel = 'rawDataCollector' +from Configuration.ProcessModifiers.gpu_cff import gpu +(~gpu).toModify(siPixelDigis, InputLabel = 'rawDataCollector') #false by default anyways ecalDigis.DoRegional = False ecalDigis.InputLabel = 'rawDataCollector' ecalPreshowerDigis.sourceTag = 'rawDataCollector' diff --git a/DQM/Integration/python/clients/beampixel_dqm_sourceclient-live_cfg.py b/DQM/Integration/python/clients/beampixel_dqm_sourceclient-live_cfg.py index b5c92413c1cf5..422e28c40a4bf 100644 --- a/DQM/Integration/python/clients/beampixel_dqm_sourceclient-live_cfg.py +++ b/DQM/Integration/python/clients/beampixel_dqm_sourceclient-live_cfg.py @@ -63,51 +63,32 @@ #---------------------------- -# Tracking Configuration -#---------------------------- -process.castorDigis.InputLabel = cms.InputTag("rawDataRepacker") -process.csctfDigis.producer = cms.InputTag("rawDataRepacker") -process.dttfDigis.DTTF_FED_Source = cms.InputTag("rawDataRepacker") -process.ecalDigis.InputLabel = cms.InputTag("rawDataRepacker") -process.ecalPreshowerDigis.sourceTag = cms.InputTag("rawDataRepacker") -process.gctDigis.inputLabel = cms.InputTag("rawDataRepacker") -process.gtDigis.DaqGtInputTag = cms.InputTag("rawDataRepacker") -process.hcalDigis.InputLabel = cms.InputTag("rawDataRepacker") -process.muonCSCDigis.InputObjects = cms.InputTag("rawDataRepacker") -process.muonDTDigis.inputLabel = cms.InputTag("rawDataRepacker") -process.muonRPCDigis.InputLabel = cms.InputTag("rawDataRepacker") -process.scalersRawToDigi.scalersInputTag = cms.InputTag("rawDataRepacker") -process.siPixelDigis.InputLabel = cms.InputTag("rawDataRepacker") -process.siStripDigis.ProductLabel = cms.InputTag("rawDataRepacker") - -process.load("RecoVertex.BeamSpotProducer.BeamSpot_cfi") -process.load("RecoLocalTracker.Configuration.RecoLocalTracker_cff") -process.load("TrackingTools.TransientTrack.TransientTrackBuilder_cfi") - - -#---------------------------- -# Pixel-Tracks&Vertices Config -#---------------------------- -from RecoPixelVertexing.PixelLowPtUtilities.siPixelClusterShapeCache_cfi import * -process.siPixelClusterShapeCachePreSplitting = siPixelClusterShapeCache.clone(src = 'siPixelClustersPreSplitting') -process.load("RecoLocalTracker.SiPixelRecHits.PixelCPEGeneric_cfi") -process.load("RecoPixelVertexing.Configuration.RecoPixelVertexing_cff") -process.pixelVertices.TkFilterParameters.minPt = process.pixelTracksTrackingRegions.RegionPSet.ptMin -process.pixelTracksTrackingRegions.RegionPSet.originRadius = 0.4 -process.pixelTracksTrackingRegions.RegionPSet.originHalfLength = 15. -process.pixelTracksTrackingRegions.RegionPSet.originXPos = 0.08 -process.pixelTracksTrackingRegions.RegionPSet.originYPos = -0.03 -process.pixelTracksTrackingRegions.RegionPSet.originZPos = 0. - - -#---------------------------- -# Proton-Proton Specific Section +# Proton-Proton Specific Part #---------------------------- if (process.runType.getRunType() == process.runType.pp_run or process.runType.getRunType() == process.runType.pp_run_stage1 or process.runType.getRunType() == process.runType.cosmic_run or process.runType.getRunType() == process.runType.cosmic_run_stage1 or process.runType.getRunType() == process.runType.hpu_run): print "[beampixel_dqm_sourceclient-live_cfg]::running pp" + process.castorDigis.InputLabel = cms.InputTag("rawDataCollector") + process.csctfDigis.producer = cms.InputTag("rawDataCollector") + process.dttfDigis.DTTF_FED_Source = cms.InputTag("rawDataCollector") + process.ecalDigis.InputLabel = cms.InputTag("rawDataCollector") + process.ecalPreshowerDigis.sourceTag = cms.InputTag("rawDataCollector") + process.gctDigis.inputLabel = cms.InputTag("rawDataCollector") + process.gtDigis.DaqGtInputTag = cms.InputTag("rawDataCollector") + process.hcalDigis.InputLabel = cms.InputTag("rawDataCollector") + process.muonCSCDigis.InputObjects = cms.InputTag("rawDataCollector") + process.muonDTDigis.inputLabel = cms.InputTag("rawDataCollector") + process.muonRPCDigis.InputLabel = cms.InputTag("rawDataCollector") + process.scalersRawToDigi.scalersInputTag = cms.InputTag("rawDataCollector") + process.siPixelDigis.InputLabel = cms.InputTag("rawDataCollector") + process.siStripDigis.ProductLabel = cms.InputTag("rawDataCollector") + + process.load("RecoVertex.BeamSpotProducer.BeamSpot_cfi") + process.load("RecoLocalTracker.Configuration.RecoLocalTracker_cff") + process.load("TrackingTools.TransientTrack.TransientTrackBuilder_cfi") + #---------------------------- # pixelVertexDQM Config @@ -134,20 +115,80 @@ minVxDoF = cms.double(10.0), minVxWgt = cms.double(0.5), fileName = cms.string("/nfshome0/dqmdev/BeamMonitorDQM/BeamPixelResults.txt")) + if process.dqmRunConfig.type.value() is "production": + process.pixelVertexDQM.fileName = cms.string("/nfshome0/dqmpro/BeamMonitorDQM/BeamPixelResults.txt") + else: + process.pixelVertexDQM.fileName = cms.string("/nfshome0/dqmdev/BeamMonitorDQM/BeamPixelResults.txt") + print "[beampixel_dqm_sourceclient-live_cfg]::saving DIP file into " + str(process.pixelVertexDQM.fileName) + + + #---------------------------- + # Pixel-Tracks&Vertices Config + #---------------------------- + from RecoPixelVertexing.PixelLowPtUtilities.siPixelClusterShapeCache_cfi import * + process.siPixelClusterShapeCachePreSplitting = siPixelClusterShapeCache.clone(src = 'siPixelClustersPreSplitting') + process.load("RecoLocalTracker.SiPixelRecHits.PixelCPEGeneric_cfi") + process.load("RecoPixelVertexing.Configuration.RecoPixelVertexing_cff") + process.pixelVertices.TkFilterParameters.minPt = process.pixelTracksTrackingRegions.RegionPSet.ptMin + process.pixelTracksTrackingRegions.RegionPSet.originRadius = 0.4 + process.pixelTracksTrackingRegions.RegionPSet.originHalfLength = 15. + process.pixelTracksTrackingRegions.RegionPSet.originXPos = 0.08 + process.pixelTracksTrackingRegions.RegionPSet.originYPos = -0.03 + process.pixelTracksTrackingRegions.RegionPSet.originZPos = 0. + + + #---------------------------- + # Pixel-Tracks&Vertices Reco + #---------------------------- + process.reconstructionStep = cms.Sequence(process.siPixelDigis* + process.siStripDigis* + process.striptrackerlocalreco* + process.offlineBeamSpot* + process.siPixelClustersPreSplitting* + process.siPixelRecHitsPreSplitting* + process.siPixelClusterShapeCachePreSplitting* + process.recopixelvertexing) + + + #---------------------------- + # Define Path + #---------------------------- + process.p = cms.Path(process.scalersRawToDigi*process.physTrigger*process.reconstructionStep*process.pixelVertexDQM*process.dqmModules) + + #---------------------------- -# Heavy Ion Specific Section +# Heavy Ion Specific Part #---------------------------- if (process.runType.getRunType() == process.runType.hi_run): print "[beampixel_dqm_sourceclient-live_cfg]::running HI" + process.castorDigis.InputLabel = cms.InputTag("rawDataRepacker") + process.csctfDigis.producer = cms.InputTag("rawDataRepacker") + process.dttfDigis.DTTF_FED_Source = cms.InputTag("rawDataRepacker") + process.ecalDigis.InputLabel = cms.InputTag("rawDataRepacker") + process.ecalPreshowerDigis.sourceTag = cms.InputTag("rawDataRepacker") + process.gctDigis.inputLabel = cms.InputTag("rawDataRepacker") + process.gtDigis.DaqGtInputTag = cms.InputTag("rawDataRepacker") + process.hcalDigis.InputLabel = cms.InputTag("rawDataRepacker") + process.muonCSCDigis.InputObjects = cms.InputTag("rawDataRepacker") + process.muonDTDigis.inputLabel = cms.InputTag("rawDataRepacker") + process.muonRPCDigis.InputLabel = cms.InputTag("rawDataRepacker") + process.scalersRawToDigi.scalersInputTag = cms.InputTag("rawDataRepacker") + process.siPixelDigis.InputLabel = cms.InputTag("rawDataRepacker") + process.siStripDigis.ProductLabel = cms.InputTag("rawDataRepacker") + + process.load("RecoVertex.BeamSpotProducer.BeamSpot_cfi") + process.load("Configuration.StandardSequences.ReconstructionHeavyIons_cff") + process.load("RecoLocalTracker.Configuration.RecoLocalTrackerHeavyIons_cff") + #---------------------------- # pixelVertexDQM Config #---------------------------- process.pixelVertexDQM = DQMEDAnalyzer('Vx3DHLTAnalyzer', - vertexCollection = cms.untracked.InputTag("pixelVertices"), + vertexCollection = cms.untracked.InputTag("hiSelectedVertexPreSplitting"), pixelHitCollection = cms.untracked.InputTag("siPixelRecHitsPreSplitting"), debugMode = cms.bool(True), nLumiFit = cms.uint32(5), @@ -168,66 +209,57 @@ minVxDoF = cms.double(10.0), minVxWgt = cms.double(0.5), fileName = cms.string("/nfshome0/dqmdev/BeamMonitorDQM/BeamPixelResults.txt")) + if process.dqmRunConfig.type.value() is "production": + process.pixelVertexDQM.fileName = cms.string("/nfshome0/dqmpro/BeamMonitorDQM/BeamPixelResults.txt") + else: + process.pixelVertexDQM.fileName = cms.string("/nfshome0/dqmdev/BeamMonitorDQM/BeamPixelResults.txt") + print "[beampixel_dqm_sourceclient-live_cfg]::saving DIP file into " + str(process.pixelVertexDQM.fileName) #---------------------------- # Pixel-Tracks&Vertices Config #---------------------------- - from RecoVertex.PrimaryVertexProducer.TkClusParameters_cff import DA_vectParameters - offlinePrimaryVertices = cms.EDProducer( - "PrimaryVertexProducer", - verbose = cms.untracked.bool(False), - TrackLabel = cms.InputTag("generalTracks"), - beamSpotLabel = cms.InputTag("offlineBeamSpot"), - TkFilterParameters = cms.PSet( - algorithm = cms.string('filter'), - maxNormalizedChi2 = cms.double(10.0), - minPixelLayersWithHits = cms.int32(2), - minSiliconLayersWithHits = cms.int32(5), - maxD0Significance = cms.double(4.0), - minPt = cms.double(0.0), - maxEta = cms.double(2.4), - trackQuality = cms.string("any")), - TkClusParameters = DA_vectParameters, - vertexCollections = cms.VPSet( - [cms.PSet(label = cms.string(""), - algorithm = cms.string("AdaptiveVertexFitter"), - chi2cutoff = cms.double(2.5), - minNdof = cms.double(0.0), - useBeamConstraint = cms.bool(False), - maxDistanceToBeam = cms.double(1.0)), - cms.PSet(label = cms.string("WithBS"), - algorithm = cms.string('AdaptiveVertexFitter'), - chi2cutoff = cms.double(2.5), - minNdof = cms.double(2.0), - useBeamConstraint = cms.bool(True), - maxDistanceToBeam = cms.double(1.0))])) - - -#---------------------------- -# File to save beamspot info -#---------------------------- -if process.dqmRunConfig.type.value() is "production": - process.pixelVertexDQM.fileName = cms.string("/nfshome0/dqmpro/BeamMonitorDQM/BeamPixelResults.txt") -else: - process.pixelVertexDQM.fileName = cms.string("/nfshome0/dqmdev/BeamMonitorDQM/BeamPixelResults.txt") -print "[beampixel_dqm_sourceclient-live_cfg]::saving DIP file into " + str(process.pixelVertexDQM.fileName) - - -#---------------------------- -# Pixel-Tracks&Vertices Reco -#---------------------------- -process.reconstructionStep = cms.Sequence(process.siPixelDigis* - process.siStripDigis* - process.striptrackerlocalreco* - process.offlineBeamSpot* - process.siPixelClustersPreSplitting* - process.siPixelRecHitsPreSplitting* - process.siPixelClusterShapeCachePreSplitting* - process.recopixelvertexing) - - -#---------------------------- -# Define Path -#---------------------------- -process.p = cms.Path(process.scalersRawToDigi*process.physTrigger*process.reconstructionStep*process.pixelVertexDQM*process.dqmModules) + from RecoPixelVertexing.PixelLowPtUtilities.siPixelClusterShapeCache_cfi import * + from RecoPixelVertexing.PixelLowPtUtilities.siPixelClusterShapeCache_cfi import * + siPixelClusterShapeCachePreSplitting = siPixelClusterShapeCache.clone(src = 'siPixelClustersPreSplitting') + + from RecoHI.HiTracking.HIPixelVerticesPreSplitting_cff import * + process.PixelLayerTriplets.BPix.HitProducer = cms.string("siPixelRecHitsPreSplitting") + process.PixelLayerTriplets.FPix.HitProducer = cms.string("siPixelRecHitsPreSplitting") + + process.hiPixel3PrimTracksFilter = process.hiFilter.clone(VertexCollection = cms.InputTag("hiSelectedVertexPreSplitting"), + clusterShapeCacheSrc = cms.InputTag("siPixelClusterShapeCachePreSplitting")) + process.hiPixel3PrimTracks.Filter = cms.InputTag("hiPixel3PrimTracksFilter") + process.hiPixel3PrimTracks.ComponentName = cms.string("GlobalTrackingRegionWithVerticesProducer") + process.hiPixel3PrimTracks.VertexCollection = cms.InputTag("hiSelectedVertexPreSplitting") + process.hiPixel3PrimTracks.ptMin = cms.double(0.9) + process.hiPixel3PrimTracks.clusterShapeCacheSrc = cms.InputTag("siPixelClusterShapeCachePreSplitting") + process.hiPixel3ProtoTracksPreSplitting.originRadius = cms.double(0.4) + process.hiPixelAdaptiveVertexPreSplitting.vertexCollections.useBeamConstraint = cms.bool(False) + + process.hiPixel3ProtoTracksPreSplitting.RegionFactoryPSet.RegionPSet.originRadius = 0.2 # default 0.2 + process.hiPixel3ProtoTracksPreSplitting.RegionFactoryPSet.RegionPSet.fixedError = 0.5 # default 3.0 + process.hiSelectedProtoTracksPreSplitting.maxD0Significance = 100 # default 5.0 + process.hiPixelAdaptiveVertexPreSplitting.TkFilterParameters.maxD0Significance = 100 # default 3.0 + process.hiPixelAdaptiveVertexPreSplitting.vertexCollections.useBeamConstraint = False # default False + process.hiPixelAdaptiveVertexPreSplitting.vertexCollections.maxDistanceToBeam = 1.0 # default 0.1 + + + #---------------------------- + # Pixel-Tracks&Vertices Reco + #---------------------------- + process.reconstructionStep = cms.Sequence(process.siPixelDigis* + process.offlineBeamSpot* + process.pixeltrackerlocalreco* + process.siPixelClusterShapeCachePreSplitting* + process.hiPixelVerticesPreSplitting* + process.PixelLayerTriplets* + process.pixelFitterByHelixProjections* + process.hiPixel3PrimTracksFilter* + process.hiPixel3PrimTracks) + + + #---------------------------- + # Define Path + #---------------------------- + process.p = cms.Path(process.scalersRawToDigi*process.physTrigger*process.reconstructionStep*process.pixelVertexDQM*process.dqmModules) diff --git a/DQM/Integration/python/clients/l1tstage2_dqm_sourceclient-live_cfg.py b/DQM/Integration/python/clients/l1tstage2_dqm_sourceclient-live_cfg.py index bb57bd794d842..01beee9fca4f5 100644 --- a/DQM/Integration/python/clients/l1tstage2_dqm_sourceclient-live_cfg.py +++ b/DQM/Integration/python/clients/l1tstage2_dqm_sourceclient-live_cfg.py @@ -114,18 +114,15 @@ # Heavy-Ion run if (process.runType.getRunType() == process.runType.hi_run): process.DQMStore.referenceFileName = "/dqmdata/dqm/reference/l1t_reference_hi.root" - process.onlineMetaDataDigis.onlineMetaDataInputLabel = cms.InputTag("rawDataRepacker") - process.onlineMetaDataRawToDigi.onlineMetaDataInputLabel = cms.InputTag("rawDataRepacker") process.castorDigis.InputLabel = cms.InputTag("rawDataRepacker") process.ctppsDiamondRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") - process.ctppsPixelDigis.inputLabel = cms.InputTag("rawDataRepacker") + process.ctppsPixelDigis.InputLabel = cms.InputTag("rawDataRepacker") process.ecalDigis.InputLabel = cms.InputTag("rawDataRepacker") process.ecalPreshowerDigis.sourceTag = cms.InputTag("rawDataRepacker") process.hcalDigis.InputLabel = cms.InputTag("rawDataRepacker") process.muonCSCDigis.InputObjects = cms.InputTag("rawDataRepacker") process.muonDTDigis.inputLabel = cms.InputTag("rawDataRepacker") process.muonRPCDigis.InputLabel = cms.InputTag("rawDataRepacker") - process.muonGEMDigis.InputLabel = cms.InputTag("rawDataRepacker") process.scalersRawToDigi.scalersInputTag = cms.InputTag("rawDataRepacker") process.siPixelDigis.InputLabel = cms.InputTag("rawDataRepacker") process.siStripDigis.ProductLabel = cms.InputTag("rawDataRepacker") @@ -133,26 +130,20 @@ process.tcdsRawToDigi.InputLabel = cms.InputTag("rawDataRepacker") process.totemRPRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") process.totemTriggerRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") - process.totemTimingRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") process.csctfDigis.producer = cms.InputTag("rawDataRepacker") process.dttfDigis.DTTF_FED_Source = cms.InputTag("rawDataRepacker") process.gctDigis.inputLabel = cms.InputTag("rawDataRepacker") process.gtDigis.DaqGtInputTag = cms.InputTag("rawDataRepacker") process.twinMuxStage2Digis.DTTM7_FED_Source = cms.InputTag("rawDataRepacker") - process.RPCTwinMuxRawToDigi.inputTag = cms.InputTag("rawDataRepacker") process.bmtfDigis.InputLabel = cms.InputTag("rawDataRepacker") - process.omtfStage2Digis.inputLabel = cms.InputTag("rawDataRepacker") process.emtfStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") process.gmtStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") - process.caloLayer1Digis.InputLabel = cms.InputTag("rawDataRepacker") process.caloStage1Digis.InputLabel = cms.InputTag("rawDataRepacker") process.caloStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") process.gtStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") process.l1tStage2CaloLayer1.fedRawDataLabel = cms.InputTag("rawDataRepacker") process.l1tStage2uGMTZeroSupp.rawData = cms.InputTag("rawDataRepacker") process.l1tStage2uGMTZeroSuppFatEvts.rawData = cms.InputTag("rawDataRepacker") - process.l1tStage2BmtfZeroSupp.rawData = cms.InputTag("rawDataRepacker") - process.l1tStage2BmtfZeroSuppFatEvts.rawData = cms.InputTag("rawDataRepacker") process.selfFatEventFilter.rawInput = cms.InputTag("rawDataRepacker") #-------------------------------------------------- diff --git a/DQM/Integration/python/clients/l1tstage2emulator_dqm_sourceclient-live_cfg.py b/DQM/Integration/python/clients/l1tstage2emulator_dqm_sourceclient-live_cfg.py index 735e5d534a9e5..0f467719e6f8d 100644 --- a/DQM/Integration/python/clients/l1tstage2emulator_dqm_sourceclient-live_cfg.py +++ b/DQM/Integration/python/clients/l1tstage2emulator_dqm_sourceclient-live_cfg.py @@ -110,18 +110,15 @@ # Heavy-Ion run if (process.runType.getRunType() == process.runType.hi_run): - process.onlineMetaDataDigis.onlineMetaDataInputLabel = cms.InputTag("rawDataRepacker") - process.onlineMetaDataRawToDigi.onlineMetaDataInputLabel = cms.InputTag("rawDataRepacker") process.castorDigis.InputLabel = cms.InputTag("rawDataRepacker") process.ctppsDiamondRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") - process.ctppsPixelDigis.inputLabel = cms.InputTag("rawDataRepacker") + process.ctppsPixelDigis.InputLabel = cms.InputTag("rawDataRepacker") process.ecalDigis.InputLabel = cms.InputTag("rawDataRepacker") process.ecalPreshowerDigis.sourceTag = cms.InputTag("rawDataRepacker") process.hcalDigis.InputLabel = cms.InputTag("rawDataRepacker") process.muonCSCDigis.InputObjects = cms.InputTag("rawDataRepacker") process.muonDTDigis.inputLabel = cms.InputTag("rawDataRepacker") process.muonRPCDigis.InputLabel = cms.InputTag("rawDataRepacker") - process.muonGEMDigis.InputLabel = cms.InputTag("rawDataRepacker") process.scalersRawToDigi.scalersInputTag = cms.InputTag("rawDataRepacker") process.siPixelDigis.InputLabel = cms.InputTag("rawDataRepacker") process.siStripDigis.ProductLabel = cms.InputTag("rawDataRepacker") @@ -129,18 +126,14 @@ process.tcdsRawToDigi.InputLabel = cms.InputTag("rawDataRepacker") process.totemRPRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") process.totemTriggerRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") - process.totemTimingRawToDigi.rawDataTag = cms.InputTag("rawDataRepacker") process.csctfDigis.producer = cms.InputTag("rawDataRepacker") process.dttfDigis.DTTF_FED_Source = cms.InputTag("rawDataRepacker") process.gctDigis.inputLabel = cms.InputTag("rawDataRepacker") process.gtDigis.DaqGtInputTag = cms.InputTag("rawDataRepacker") process.twinMuxStage2Digis.DTTM7_FED_Source = cms.InputTag("rawDataRepacker") - process.RPCTwinMuxRawToDigi.inputTag = cms.InputTag("rawDataRepacker") process.bmtfDigis.InputLabel = cms.InputTag("rawDataRepacker") - process.omtfStage2Digis.inputLabel = cms.InputTag("rawDataRepacker") process.emtfStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") process.gmtStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") - process.caloLayer1Digis.InputLabel = cms.InputTag("rawDataRepacker") process.caloStage1Digis.InputLabel = cms.InputTag("rawDataRepacker") process.caloStage2Digis.InputLabel = cms.InputTag("rawDataRepacker") process.simHcalTriggerPrimitiveDigis.InputTagFEDRaw = cms.InputTag("rawDataRepacker") diff --git a/EventFilter/SiPixelRawToDigi/python/SiPixelRawToDigi_cfi.py b/EventFilter/SiPixelRawToDigi/python/SiPixelRawToDigi_cfi.py index fda0b69fa46a1..528ffa2683fa8 100644 --- a/EventFilter/SiPixelRawToDigi/python/SiPixelRawToDigi_cfi.py +++ b/EventFilter/SiPixelRawToDigi/python/SiPixelRawToDigi_cfi.py @@ -1,5 +1,6 @@ import FWCore.ParameterSet.Config as cms import EventFilter.SiPixelRawToDigi.siPixelRawToDigi_cfi +import RecoLocalTracker.SiPixelClusterizer.siPixelDigiHeterogeneousConverter_cfi siPixelDigis = EventFilter.SiPixelRawToDigi.siPixelRawToDigi_cfi.siPixelRawToDigi.clone() siPixelDigis.Timing = cms.untracked.bool(False) @@ -20,3 +21,10 @@ from Configuration.Eras.Modifier_phase1Pixel_cff import phase1Pixel phase1Pixel.toModify(siPixelDigis, UsePhase1=True) + +_siPixelDigis_gpu = RecoLocalTracker.SiPixelClusterizer.siPixelDigiHeterogeneousConverter_cfi.siPixelDigiHeterogeneousConverter.clone() +_siPixelDigis_gpu.includeErrors = cms.bool(True) + +from Configuration.ProcessModifiers.gpu_cff import gpu +gpu.toReplaceWith(siPixelDigis, _siPixelDigis_gpu) + 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 0692d49f0068f..ad975fa183566 100644 --- a/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py +++ b/RecoLocalTracker/Configuration/python/RecoLocalTracker_cff.py @@ -17,6 +17,10 @@ striptrackerlocalreco = cms.Sequence(siStripZeroSuppression*siStripClusters*siStripMatchedRecHits) trackerlocalreco = cms.Sequence(pixeltrackerlocalreco*striptrackerlocalreco*clusterSummaryProducer) +from Configuration.ProcessModifiers.gpu_cff import gpu +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHeterogeneous_cfi import siPixelRecHitHeterogeneous as _siPixelRecHitHeterogeneous +gpu.toReplaceWith(siPixelRecHitsPreSplitting, _siPixelRecHitHeterogeneous) + from RecoLocalTracker.SiPhase2Clusterizer.phase2TrackerClusterizer_cfi import * from RecoLocalTracker.Phase2TrackerRecHits.Phase2StripCPEGeometricESProducer_cfi import * diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml b/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml index f6958151bdb91..9db4a46f367b3 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml @@ -3,7 +3,16 @@ + + - + + + + + + + + diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index f3242a11d7ae6..7bd6eac473cc7 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -34,6 +34,7 @@ #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h" +#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h" #include "RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPU.h" // local includes @@ -687,11 +688,11 @@ namespace pixelgpudetails { cudaCheck(cudaMemcpyAsync(adc_h, adc_d, wordCounter*sizeof(uint16_t), cudaMemcpyDefault, stream.id())); } - /* - std::cout +#ifdef GPU_DEBUG + std::cout << "CUDA countModules kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; - */ +#endif cudaCheck(cudaMemsetAsync(moduleStart_d, 0x00, sizeof(uint32_t), stream.id())); @@ -703,10 +704,10 @@ namespace pixelgpudetails { threadsPerBlock = 256; blocks = MaxNumModules; - /* +#ifdef GPU_DEBUG std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; - */ +#endif cudaCheck(cudaMemsetAsync(clusInModule_d, 0, (MaxNumModules)*sizeof(uint32_t), stream.id())); findClus<<>>( moduleInd_d, @@ -717,6 +718,18 @@ namespace pixelgpudetails { wordCounter); cudaCheck(cudaGetLastError()); + // apply charge cut + clusterChargeCut<<>>( + moduleInd_d, + adc_d, + moduleStart_d, + clusInModule_d, moduleId_d, + clus_d, + wordCounter); + cudaCheck(cudaGetLastError()); + + + // count the module start indices already here (instead of // rechits) so that the number of clusters/hits can be made // available in the rechit producer without additional points of diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h new file mode 100644 index 0000000000000..0284a378ecd39 --- /dev/null +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -0,0 +1,97 @@ +#ifndef RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterChargeCut_h +#define RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterChargeCut_h + +#include +#include +#include + +#include "gpuClusteringConstants.h" + +#include "HeterogeneousCore/CUDAUtilities/interface/prefixScan.h" + +namespace gpuClustering { + + __global__ void clusterChargeCut( + uint16_t * __restrict__ id, // module id of each pixel (modified if bad cluster) + uint16_t const * __restrict__ adc, // charge of each pixel + uint32_t const * __restrict__ moduleStart, // index of the first pixel of each module + uint32_t * __restrict__ nClustersInModule, // modified: number of clusters found in each module + uint32_t const * __restrict__ moduleId, // module id of each module + int32_t * __restrict__ clusterId, // modified: cluster id of each pixel + int numElements) + { + + if (blockIdx.x >= moduleStart[0]) + return; + + auto firstPixel = moduleStart[1 + blockIdx.x]; + auto thisModuleId = id[firstPixel]; + assert(thisModuleId < MaxNumModules); + assert(thisModuleId==moduleId[blockIdx.x]); + + auto nclus = nClustersInModule[thisModuleId]; + if (nclus==0) return; + + assert(nclus<=MaxNumClustersPerModules); + +#ifdef GPU_DEBUG + if (thisModuleId % 100 == 1) + if (threadIdx.x == 0) + printf("start clusterizer for module %d in block %d\n", thisModuleId, blockIdx.x); +#endif + + auto first = firstPixel + threadIdx.x; + + __shared__ int32_t charge[MaxNumClustersPerModules]; + for (int i=threadIdx.x; ichargeCut ? 1 : 0; + } + + __syncthreads(); + + // renumber + __shared__ uint16_t ws[32]; + blockPrefixScan(newclusId, nclus, ws); + + assert(nclus>=newclusId[nclus-1]); + + if(nclus==newclusId[nclus-1]) return; + + nClustersInModule[thisModuleId] = newclusId[nclus-1]; + __syncthreads(); + + // mark bad cluster again + for (int i=threadIdx.x; i #include #include #include "gpuClusteringConstants.h" +#include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h" #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h" + namespace gpuClustering { __global__ void countModules(uint16_t const * __restrict__ id, @@ -32,7 +37,9 @@ namespace gpuClustering { } } - __global__ void findClus(uint16_t const * __restrict__ id, // module id of each pixel + __global__ +// __launch_bounds__(256,4) + void findClus(uint16_t const * __restrict__ id, // module id of each pixel uint16_t const * __restrict__ x, // local coordinates of each pixel uint16_t const * __restrict__ y, // uint32_t const * __restrict__ moduleStart, // index of the first pixel of each module @@ -63,8 +70,6 @@ namespace gpuClustering { __syncthreads(); // skip threads not associated to an existing pixel - bool active = (first < numElements); - if (active) { for (int i = first; i < numElements; i += blockDim.x) { if (id[i] == InvId) // skip invalid pixels continue; @@ -73,95 +78,132 @@ namespace gpuClustering { break; } } - } - - //init hist (ymax < 512) - __shared__ HistoContainer hist; - hist.nspills = 0; - for (auto k = threadIdx.x; k; + constexpr auto wss = Hist::totbins(); + __shared__ Hist hist; + __shared__ typename Hist::Counter ws[wss]; + for (auto j=threadIdx.x; j0 ? y[i]-1 : 0); - auto be = hist.bin(y[i]+1)+1; - auto loop = [&](int j) { - j+=firstPixel; - if (i>=j or j>jm or - std::abs(int(x[j]) - int(x[i])) > 1 or - std::abs(int(y[j]) - int(y[i])) > 1) return; - auto old = atomicMin(&clusterId[j], clusterId[i]); + 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]); if (old != clusterId[i]) { // end the loop only if no changes were applied - done = false; + more = true; } atomicMin(&clusterId[i], old); +#ifdef CLUS_LIMIT_LOOP // update the loop boundary for the next iteration - jmax[k] = std::max(j + 1,jmax[k]); + jmax[k] = std::max(kk + 1,jmax[k]); +#endif }; - for (auto b=bs; b MaxNumModules + } #endif // RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusteringConstants_h diff --git a/RecoLocalTracker/SiPixelClusterizer/python/SiPixelClusterizerPreSplitting_cfi.py b/RecoLocalTracker/SiPixelClusterizer/python/SiPixelClusterizerPreSplitting_cfi.py index ba8d492c5f610..bb0bb85697a99 100644 --- a/RecoLocalTracker/SiPixelClusterizer/python/SiPixelClusterizerPreSplitting_cfi.py +++ b/RecoLocalTracker/SiPixelClusterizer/python/SiPixelClusterizerPreSplitting_cfi.py @@ -1,7 +1,11 @@ - import FWCore.ParameterSet.Config as cms -# from CondTools.SiPixel.SiPixelGainCalibrationService_cfi import * from RecoLocalTracker.SiPixelClusterizer.SiPixelClusterizer_cfi import siPixelClusters as _siPixelClusters siPixelClustersPreSplitting = _siPixelClusters.clone() + +from Configuration.ProcessModifiers.gpu_cff import gpu +from RecoLocalTracker.SiPixelClusterizer.siPixelClustersHeterogeneous_cfi import siPixelClustersHeterogeneous as _siPixelClustersHeterogeneous +from RecoLocalTracker.SiPixelClusterizer.siPixelFedCablingMapGPUWrapper_cfi import * +from CalibTracker.SiPixelESProducers.siPixelGainCalibrationForHLTGPU_cfi import * +gpu.toReplaceWith(siPixelClustersPreSplitting, _siPixelClustersHeterogeneous.clone()) diff --git a/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml b/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml index 3445783781551..335591b583b58 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml +++ b/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml @@ -31,3 +31,18 @@ + + + + + + + + + + + + + + + diff --git a/RecoLocalTracker/SiPixelRecHits/BuildFile.xml b/RecoLocalTracker/SiPixelRecHits/BuildFile.xml index 4fc33a07b7477..7918c7a4f4d9a 100644 --- a/RecoLocalTracker/SiPixelRecHits/BuildFile.xml +++ b/RecoLocalTracker/SiPixelRecHits/BuildFile.xml @@ -12,6 +12,11 @@ + + + + + diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml b/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml index 0b477d0fb2d34..a8af0c8a7c4f9 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml +++ b/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml @@ -1,7 +1,12 @@ - - - - - - + + + + + + + + + + + diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu index e2165471c3386..c63466f157a1b 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu @@ -8,7 +8,7 @@ // CMSSW headers #include "RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" -#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h" +#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h" #include "PixelRecHits.h" #include "gpuPixelRecHits.h" @@ -34,7 +34,7 @@ namespace { namespace pixelgpudetails { PixelRecHitGPUKernel::PixelRecHitGPUKernel(cuda::stream_t<>& cudaStream) { - constexpr auto MAX_HITS = gpuClustering::MaxNumModules * 256; + constexpr auto MAX_HITS = siPixelRecHitsHeterogeneousProduct::maxHits(); cudaCheck(cudaMalloc((void **) & gpu_.bs_d, 3 * sizeof(float))); cudaCheck(cudaMalloc((void **) & gpu_.hitsLayerStart_d, 11 * sizeof(uint32_t))); @@ -70,7 +70,8 @@ namespace pixelgpudetails { gpu_.iphi_d = slicePitch(gpu_.owner_16bit_, gpu_.owner_16bit_pitch_, 3); gpu_.sortIndex_d = slicePitch(gpu_.owner_16bit_, gpu_.owner_16bit_pitch_, 4); - cudaCheck(cudaMalloc((void **) & gpu_.hist_d, 10 * sizeof(HitsOnGPU::Hist))); + cudaCheck(cudaMalloc((void **) & gpu_.hist_d, sizeof(HitsOnGPU::Hist))); + cudaCheck(cudaMalloc((void **) & gpu_.hws_d, 4*HitsOnGPU::Hist::totbins())); cudaCheck(cudaMalloc((void **) & gpu_d, sizeof(HitsOnGPU))); gpu_.me_d = gpu_d; cudaCheck(cudaMemcpyAsync(gpu_d, &gpu_, sizeof(HitsOnGPU), cudaMemcpyDefault, cudaStream.id())); @@ -109,6 +110,7 @@ namespace pixelgpudetails { cudaCheck(cudaFree(gpu_.owner_32bit_)); cudaCheck(cudaFree(gpu_.owner_16bit_)); cudaCheck(cudaFree(gpu_.hist_d)); + cudaCheck(cudaFree(gpu_.hws_d)); cudaCheck(cudaFree(gpu_d)); cudaCheck(cudaFree(d_phase1TopologyLayerStart_)); @@ -132,6 +134,10 @@ namespace pixelgpudetails { int threadsPerBlock = 256; int blocks = input.nModules; // active modules (with digis) + +#ifdef GPU_DEBUG + std::cout << "launching getHits kernel for " << blocks << " blocks" << std::endl; +#endif gpuPixelRecHits::getHits<<>>( cpeParams, gpu_.bs_d, @@ -188,6 +194,6 @@ namespace pixelgpudetails { // radixSortMultiWrapper<<<10, 256, 0, c.stream>>>(gpu_.iphi_d, gpu_.sortIndex_d, gpu_.hitsLayerStart_d); } - cudautils::fillManyFromVector(gpu_.hist_d, 10, gpu_.iphi_d, gpu_.hitsLayerStart_d, nhits_, 256, stream.id()); + cudautils::fillManyFromVector(gpu_.hist_d, gpu_.hws_d, 10, gpu_.iphi_d, gpu_.hitsLayerStart_d, nhits_, 256, stream.id()); } } diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index 2ee4a10c6fc99..3f92e4833bc22 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -14,15 +14,6 @@ namespace gpuPixelRecHits { - // to be moved in common namespace... - constexpr uint16_t InvId=9999; // must be > MaxNumModules - - - constexpr uint32_t MaxClusInModule = pixelCPEforGPU::MaxClusInModule; - - using ClusParams = pixelCPEforGPU::ClusParams; - - __global__ void getHits(pixelCPEforGPU::ParamsOnGPU const * __restrict__ cpeParams, float const * __restrict__ bs, uint16_t const * __restrict__ id, @@ -42,21 +33,44 @@ namespace gpuPixelRecHits { float * xe, float * ye, uint16_t * mr, uint16_t * mc) { + + // to be moved in common namespace... + constexpr uint16_t InvId=9999; // must be > MaxNumModules + constexpr uint32_t MaxClusInModule = pixelCPEforGPU::MaxClusInModule; + + using ClusParams = pixelCPEforGPU::ClusParams; + + // as usual one block per module __shared__ ClusParams clusParams; auto first = digiModuleStart[1 + blockIdx.x]; - auto me = id[first]; - assert(moduleId[blockIdx.x] == me); + auto me = moduleId[blockIdx.x]; auto nclus = clusInModule[me]; + if (0==nclus) return; + +#ifdef GPU_DEBUG + if (threadIdx.x==0) { + auto k=first; + while (id[k]==InvId) ++k; + assert(id[k]==me); + } +#endif + #ifdef GPU_DEBUG if (me%100==1) if (threadIdx.x==0) printf("hitbuilder: %d clusters in module %d. will write at %d\n", nclus, me, hitsModuleStart[me]); #endif assert(blockDim.x >= MaxClusInModule); - assert(nclus <= MaxClusInModule); + + if (threadIdx.x==0 && nclus > MaxClusInModule) { + printf("WARNING: too many clusters %d in Module %d. Only first %d processed\n", nclus,me,MaxClusInModule); + // zero charge: do not bother to do it in parallel + for (auto d=MaxClusInModule; d= nclus) continue; atomicMin(&clusParams.minRow[clus[i]], x[i]); atomicMax(&clusParams.maxRow[clus[i]], x[i]); atomicMin(&clusParams.minCol[clus[i]], y[i]); @@ -93,6 +107,7 @@ namespace gpuPixelRecHits { for (int i = first; i < numElements; i += blockDim.x) { if (id[i] == InvId) continue; // not valid if (id[i] != me) break; // end of module + if (clus[i] >= nclus) continue; atomicAdd(&clusParams.charge[clus[i]], adc[i]); if (clusParams.minRow[clus[i]]==x[i]) atomicAdd(&clusParams.Q_f_X[clus[i]], adc[i]); if (clusParams.maxRow[clus[i]]==x[i]) atomicAdd(&clusParams.Q_l_X[clus[i]], adc[i]); diff --git a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py index a1ff25af2e697..8e28bbb175181 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py +++ b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py @@ -19,6 +19,7 @@ # 4. Pixel Generic CPE # from RecoLocalTracker.SiPixelRecHits.PixelCPEGeneric_cfi import * +from RecoLocalTracker.SiPixelRecHits.PixelCPEFast_cfi import * # # 5. ESProducer for the Magnetic-field dependent template records #