From 1e813470368a11cefce5fcb8fabfe7c0db0f636f Mon Sep 17 00:00:00 2001 From: AdrianoDee Date: Fri, 30 Jun 2023 14:51:00 +0200 Subject: [PATCH 1/2] Pixel tracks updated: - HIon setup - More configurable chain --- .../interface/gpuClusteringConstants.h | 4 +- .../interface/TrackSoAHeterogeneousDevice.h | 1 + .../interface/TrackSoAHeterogeneousHost.h | 2 +- CUDADataFormats/Track/src/classes_def.xml | 6 + .../interface/TrackingRecHitSoADevice.h | 2 + .../interface/TrackingRecHitSoAHost.h | 2 + .../TrackingRecHit/src/classes_def.xml | 7 +- .../PyReleaseValidation/python/relval_gpu.py | 2 + .../python/relval_standard.py | 2 + .../python/relval_steps.py | 12 ++ .../python/upgradeWorkflowComponents.py | 79 ++++++- .../plugins/SiPixelCompareRecHitsSoA.cc | 3 + .../plugins/SiPixelCompareTrackSoA.cc | 3 + .../plugins/SiPixelMonitorRecHitsSoA.cc | 3 + .../plugins/SiPixelMonitorTrackSoA.cc | 2 + .../SiPixelHeterogenousDQM_FirstStep_cff.py | 41 ++++ .../interface/SimplePixelTopology.h | 64 +++++- .../python/HILowPtConformalPixelTracks_cfi.py | 85 +++++++- .../plugins/SiPixelClusterThresholds.h | 49 ++++- .../plugins/SiPixelDigisClustersFromSoA.cc | 14 +- .../plugins/SiPixelPhase2DigiToClusterCUDA.cc | 19 +- .../plugins/SiPixelRawToClusterCUDA.cc | 74 ++++--- .../plugins/SiPixelRawToClusterGPUKernel.cu | 199 +++++++++--------- .../plugins/SiPixelRawToClusterGPUKernel.h | 34 +-- .../plugins/gpuCalibPixel.h | 52 ++--- .../plugins/gpuClusterChargeCut.h | 24 ++- .../plugins/gpuClustering.h | 3 + .../python/siPixelClustersPreSplitting_cff.py | 29 ++- .../SiPixelClusterizer/test/gpuClustering_t.h | 2 +- .../plugins/PixelCPEFastESProducer.cc | 2 + .../plugins/PixelRecHitGPUKernel.cu | 1 + .../plugins/SiPixelRecHitCUDA.cc | 3 + .../plugins/SiPixelRecHitFromCUDA.cc | 5 +- .../plugins/SiPixelRecHitSoAFromCUDA.cc | 3 + .../plugins/SiPixelRecHitSoAFromLegacy.cc | 3 + .../python/PixelCPEESProducers_cff.py | 1 + .../python/SiPixelRecHits_cfi.py | 21 +- .../SiPixelRecHits/src/PixelCPEFast.cc | 1 + .../python/RecoPixelVertexing_cff.py | 16 ++ .../plugins/BrokenLineFitOnGPU.cc | 1 + .../plugins/BrokenLineFitOnGPU.cu | 1 + .../PixelSeeding/plugins/CAHitNtupletCUDA.cc | 3 + .../plugins/CAHitNtupletGeneratorKernels.cc | 9 +- .../plugins/CAHitNtupletGeneratorKernels.cu | 13 +- .../plugins/CAHitNtupletGeneratorKernels.h | 11 +- .../CAHitNtupletGeneratorKernelsAlloc.cc | 11 +- .../CAHitNtupletGeneratorKernelsImpl.h | 14 +- .../plugins/CAHitNtupletGeneratorOnGPU.cc | 80 +++++-- .../PixelSeeding/plugins/HelixFitOnGPU.cc | 1 + .../PixelSeeding/plugins/RiemannFitOnGPU.cc | 1 + .../PixelSeeding/plugins/RiemannFitOnGPU.cu | 1 + .../PixelSeeding/plugins/gpuPixelDoublets.h | 5 +- .../plugins/gpuPixelDoubletsAlgos.h | 64 ++++-- RecoTracker/PixelSeeding/test/trip_cfg.py | 4 +- .../plugins/PixelTrackDumpCUDA.cc | 3 + .../plugins/PixelTrackProducerFromSoA.cc | 3 + .../plugins/PixelTrackSoAFromCUDA.cc | 3 + .../python/PixelTracks_cff.py | 23 ++ .../plugins/PixelVertexProducerCUDA.cc | 5 + .../plugins/gpuVertexFinder.cc | 19 +- .../plugins/gpuVertexFinder.h | 3 + .../python/PostProcessorTracker_cfi.py | 16 +- .../RecoTrack/python/TrackValidation_cff.py | 58 ++++- .../python/plotting/trackingPlots.py | 2 +- 64 files changed, 946 insertions(+), 283 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index 05b59c56559d1..cb2b7ace8bd17 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -13,9 +13,9 @@ namespace gpuClustering { // tested on MC events with 55-75 pileup events constexpr uint32_t maxHitsInIter() { return 160; } //TODO better tuning for PU 140-200 #endif - constexpr uint32_t maxHitsInModule() { return 1024; } + constexpr uint32_t maxHitsInModule() { return 2048; } - constexpr uint32_t maxNumDigis = 3 * 256 * 1024; // @PU=200 µ=530k σ=50k this is >4σ away + constexpr uint32_t maxNumDigis = 3 * 256 * 1024; // @PU=200 µ=530 sigma=50k this is >4sigma away constexpr uint16_t maxNumModules = 4000; constexpr int32_t maxNumClustersPerModules = maxHitsInModule(); diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h index 1938991e071e1..04d286a767ab0 100644 --- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h +++ b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousDevice.h @@ -30,6 +30,7 @@ namespace pixelTrack { using TrackSoADevicePhase1 = TrackSoAHeterogeneousDevice; using TrackSoADevicePhase2 = TrackSoAHeterogeneousDevice; + using TrackSoADeviceHIonPhase1 = TrackSoAHeterogeneousDevice; } // namespace pixelTrack diff --git a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h index af8af2a40a52e..39e83491e1769 100644 --- a/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h +++ b/CUDADataFormats/Track/interface/TrackSoAHeterogeneousHost.h @@ -29,7 +29,7 @@ namespace pixelTrack { using TrackSoAHostPhase1 = TrackSoAHeterogeneousHost; using TrackSoAHostPhase2 = TrackSoAHeterogeneousHost; - + using TrackSoAHostHIonPhase1 = TrackSoAHeterogeneousHost; } // namespace pixelTrack #endif // CUDADataFormats_Track_TrackHeterogeneousT_H diff --git a/CUDADataFormats/Track/src/classes_def.xml b/CUDADataFormats/Track/src/classes_def.xml index 5e3116609330a..5314f3f20b0d7 100644 --- a/CUDADataFormats/Track/src/classes_def.xml +++ b/CUDADataFormats/Track/src/classes_def.xml @@ -12,4 +12,10 @@ + + + + + + diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index ab0043930558d..0a585792ca158 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h @@ -77,5 +77,7 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection; using TrackingRecHitSoADevicePhase2 = TrackingRecHitSoADevice; +using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice; +using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice; #endif // CUDADataFormats_Track_TrackHeterogeneousT_H diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h index fdbf1432a6442..6382645b7cb5b 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h @@ -67,5 +67,7 @@ class TrackingRecHitSoAHost : public cms::cuda::PortableHostCollection; using TrackingRecHitSoAHostPhase2 = TrackingRecHitSoAHost; +using TrackingRecHitSoAHostHIonPhase1 = TrackingRecHitSoAHost; +using TrackingRecHitSoAHostHIonPhase1 = TrackingRecHitSoAHost; #endif // CUDADataFormats_Track_TrackHeterogeneousT_H diff --git a/CUDADataFormats/TrackingRecHit/src/classes_def.xml b/CUDADataFormats/TrackingRecHit/src/classes_def.xml index 6c2389e829549..dfc2c6d748e0f 100644 --- a/CUDADataFormats/TrackingRecHit/src/classes_def.xml +++ b/CUDADataFormats/TrackingRecHit/src/classes_def.xml @@ -3,14 +3,19 @@ + + + - + + + diff --git a/Configuration/PyReleaseValidation/python/relval_gpu.py b/Configuration/PyReleaseValidation/python/relval_gpu.py index 6374cf9e65267..78a82711cc678 100644 --- a/Configuration/PyReleaseValidation/python/relval_gpu.py +++ b/Configuration/PyReleaseValidation/python/relval_gpu.py @@ -71,3 +71,5 @@ workflows[141.008506] = ['Run3-2023_JetMET2023B_RecoPixelOnlyTripletsGPU',['RunJetMET2023B','HLTDR3_2023','RECODR3_reHLT_Patatrack_PixelOnlyTripletsGPU','HARVESTRUN3_pixelTrackingOnly']] workflows[141.008512] = ['Run3-2023_JetMET2023B_RecoECALOnlyGPU',['RunJetMET2023B','HLTDR3_2023','RECODR3_reHLT_ECALOnlyGPU','HARVESTRUN3_ECALOnly']] workflows[141.008522] = ['Run3-2023_JetMET2023B_RecoHCALOnlyGPU',['RunJetMET2023B','HLTDR3_2023','RECODR3_reHLT_HCALOnlyGPU','HARVESTRUN3_HCALOnly']] + +workflows[160.502] = ['',['HydjetQ_MinBias_5362GeV_2023_ppReco','DIGIHI2023PPRECO','RAWPRIMESIMHI18','RECOHI2023PPRECOMB_PatatrackGPU','MINIHI2023PROD']] \ No newline at end of file diff --git a/Configuration/PyReleaseValidation/python/relval_standard.py b/Configuration/PyReleaseValidation/python/relval_standard.py index 9a8aab2133d03..4f1f43445a417 100644 --- a/Configuration/PyReleaseValidation/python/relval_standard.py +++ b/Configuration/PyReleaseValidation/python/relval_standard.py @@ -824,6 +824,8 @@ workflows[160.2] = ['',['PhotonJets_Pt_10_5362_HI_2023','DIGIHI2023PPRECO','RECOHI2023PPRECO','HARVESTHI2023PPRECO']] workflows[160.3] = ['',['ZMM_5362_HI_2023','DIGIHI2023PPRECO','RECOHI2023PPRECO','HARVESTHI2023PPRECO']] workflows[160.4] = ['',['ZEE_5362_HI_2023','DIGIHI2023PPRECO','RECOHI2023PPRECO','HARVESTHI2023PPRECO']] +# Patatrack Pixel Tracks on CPU +workflows[160.501] = ['',['HydjetQ_MinBias_5362GeV_2023_ppReco','DIGIHI2023PPRECO','RAWPRIMESIMHI18','RECOHI2023PPRECOMB_PatatrackCPU','MINIHI2023PROD']] ### pp reference test ### workflows[149] = ['',['QCD_Pt_80_120_13_PPREF','DIGIPPREF2017','RECOPPREF2017','HARVESTPPREF2017']] diff --git a/Configuration/PyReleaseValidation/python/relval_steps.py b/Configuration/PyReleaseValidation/python/relval_steps.py index 15ce19ed725f4..1df1d0d35e29b 100644 --- a/Configuration/PyReleaseValidation/python/relval_steps.py +++ b/Configuration/PyReleaseValidation/python/relval_steps.py @@ -3013,6 +3013,18 @@ def gen2023HiMix(fragment,howMuch): '--era':'Run3_pp_on_PbPb_2023', '--procModifiers':'genJetSubEvent', },step3Up2015Defaults]) +steps['RECOHI2023PPRECOMB_PatatrackGPU']=merge([hiDefaults2023_ppReco,step3_pixel_ntuplet_gpu,{'-s':'RAW2DIGI,L1Reco,RECO,PAT,VALIDATION:@standardValidationNoHLT+@miniAODValidation,DQM:@standardDQMFakeHLT+@miniAODDQM', + '--datatier':'GEN-SIM-RECO,MINIAODSIM,DQMIO', + '--eventcontent':'RECOSIM,MINIAODSIM,DQM', + '--era':'Run3_pp_on_PbPb', + '--procModifiers':'genJetSubEvent', + },step3Up2015Defaults]) +steps['RECOHI2023PPRECOMB_PatatrackCPU']=merge([hiDefaults2023_ppReco,step3_pixel_ntuplet_cpu,{'-s':'RAW2DIGI,L1Reco,RECO,PAT,VALIDATION:@standardValidationNoHLT+@miniAODValidation,DQM:@standardDQMFakeHLT+@miniAODDQM', + '--datatier':'GEN-SIM-RECO,MINIAODSIM,DQMIO', + '--eventcontent':'RECOSIM,MINIAODSIM,DQM', + '--era':'Run3_pp_on_PbPb', + '--procModifiers':'genJetSubEvent', + },step3Up2015Defaults]) steps['REMINIAODHI2023PPRECOMB']=merge([{'-s':'PAT,VALIDATION:@miniAODValidation,DQM:@miniAODDQM', '--datatier':'MINIAODSIM,DQMIO', '--eventcontent':'MINIAODSIM,DQM', diff --git a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py index ddc2dc23b26ef..0a00bfcd98b73 100644 --- a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py +++ b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py @@ -30,6 +30,10 @@ '2021postEEPU', '2023FS', '2023FSPU', + '2022HI', + '2022HIRP', #RawPrime + '2023HI', + '2023HIRP', #RawPrime ] upgradeKeys[2026] = [ @@ -249,7 +253,7 @@ def __init__(self, steps, PU, suffix, offset): steps = steps + ["ALCA","Nano"] super().__init__(steps, PU, suffix, offset) def condition(self, fragment, stepList, key, hasHarvest): - result = (fragment=="TTbar_13" or fragment=="TTbar_14TeV") and not 'PU' in key and hasHarvest and self.condition_(fragment, stepList, key, hasHarvest) + result = (fragment=="TTbar_13" or fragment=="TTbar_14TeV" or 'Hydjet' in fragment) and not 'PU' in key and hasHarvest and self.condition_(fragment, stepList, key, hasHarvest) return result def condition_(self, fragment, stepList, key, hasHarvest): return True @@ -370,7 +374,7 @@ def setup__(self, step, stepName, stepDict, k, properties): elif 'ALCA' in step: stepDict[stepName][k] = None elif 'HARVEST' in step: stepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, stepDict[step][k]]) def condition_(self, fragment, stepList, key, hasHarvest): - return ('2017' in key or '2018' in key or '2021' in key or '2026' in key) and ('FS' not in key) + return ('2017' in key or '2018' in key or '2021' in key or '2026' in key or 'HI' in key) and ('FS' not in key) upgradeWFs['pixelTrackingOnly'] = UpgradeWorkflow_pixelTrackingOnly( steps = [ 'Reco', @@ -831,7 +835,8 @@ def condition(self, fragment, stepList, key, hasHarvest): ('2018' in key and fragment == "ZMM_13"), ('2021' in key and fragment == "ZMM_14" and 'FS' not in key), ('2023' in key and fragment == "ZMM_14" and 'FS' not in key), - ('2026D88' in key and fragment == "TTbar_14TeV" and "PixelOnly" in self.suffix) + ('2026D88' in key and fragment == "TTbar_14TeV" and "PixelOnly" in self.suffix), + (('HI' in key) and 'Hydjet' in fragment and "PixelOnly" in self.suffix ) ] result = any(selected) and hasHarvest @@ -1485,6 +1490,37 @@ def setup_(self, step, stepName, stepDict, k, properties): offset = 0.597, ) + +class PatatrackWorkflowHI(PatatrackWorkflow): + + def condition(self, fragment, stepList, key, hasHarvest): + # select only a subset of the workflows + selected = [ + ('Hydjet_Quenched' in fragment and "PixelOnly" in self.suffix ) + ] + result = any(selected) and hasHarvest + + return result + + def setup_(self, step, stepName, stepDict, k, properties): + # skip ALCA and Nano steps (but not RecoNano or HARVESTNano for Run3) + if 'ALCA' in step or 'Nano'==step: + stepDict[stepName][k] = None + elif 'Digi' in step: + if self.__digi is None: + stepDict[stepName][k] = None + else: + stepDict[stepName][k] = merge([self.__digi, stepDict[step][k]]) + elif 'Reco' in step: + if self.__reco is None: + stepDict[stepName][k] = None + else: + stepDict[stepName][k] = merge([self.__reco, stepDict[step][k]]) + elif 'HARVEST' in step: + if self.__harvest is None: + stepDict[stepName][k] = None + else: + stepDict[stepName][k] = merge([self.__harvest, stepDict[step][k]]) # end of Patatrack workflows class UpgradeWorkflow_ProdLike(UpgradeWorkflow): @@ -2403,7 +2439,8 @@ def condition(self, fragment, stepList, key, hasHarvest): class UpgradeWorkflow_DDDDB(UpgradeWorkflow): def setup_(self, step, stepName, stepDict, k, properties): - if 'Run3' in stepDict[step][k]['--era'] and '2023' not in stepDict[step][k]['--era'] and 'Fast' not in stepDict[step][k]['--era']: + theEra = stepDict[step][k]['--era'] + if 'Run3' in stepDict[step][k]['--era'] and '2023' not in stepDict[step][k]['--era'] and 'Fast' not in theEra and "Pb" not in theEra: # retain any other eras tmp_eras = stepDict[step][k]['--era'].split(',') tmp_eras[tmp_eras.index("Run3")] = 'Run3_DDD' @@ -2580,6 +2617,38 @@ def condition(self, fragment, stepList, key, hasHarvest): 'BeamSpot': 'Realistic25ns13p6TeVEarly2022Collision', 'ScenToRun' : ['Gen','FastSimRun3','HARVESTFastRun3'], }, + '2022HI' : { + 'Geom' : 'DB:Extended', + 'GT':'auto:phase1_2022_realistic_hi', + 'HLTmenu': '@fake2', + 'Era':'Run3_pp_on_PbPb', + 'BeamSpot': 'Realistic2022PbPbCollision', + 'ScenToRun' : ['GenSim','Digi','RecoNano','HARVESTNano','ALCA'], + }, + '2022HIRP' : { + 'Geom' : 'DB:Extended', + 'GT':'auto:phase1_2022_realistic_hi', + 'HLTmenu': '@fake2', + 'Era':'Run3_pp_on_PbPb_approxSiStripClusters', + 'BeamSpot': 'Realistic2022PbPbCollision', + 'ScenToRun' : ['GenSim','Digi','RecoNano','HARVESTNano','ALCA'], + }, + '2023HI' : { + 'Geom' : 'DB:Extended', + 'GT':'auto:phase1_2023_realistic_hi', + 'HLTmenu': '@fake2', + 'Era':'Run3_pp_on_PbPb', + 'BeamSpot': 'Realistic2022PbPbCollision', + 'ScenToRun' : ['GenSim','Digi','RecoNano','HARVESTNano','ALCA'], + }, + '2023HIRP' : { + 'Geom' : 'DB:Extended', + 'GT':'auto:phase1_2023_realistic_hi', + 'HLTmenu': '@fake2', + 'Era':'Run3_pp_on_PbPb_approxSiStripClusters', + 'BeamSpot': 'Realistic2022PbPbCollision', + 'ScenToRun' : ['GenSim','Digi','RecoNano','HARVESTNano','ALCA'], + } } # standard PU sequences @@ -2858,4 +2927,6 @@ def __init__(self, howMuch, dataset): ('LbToJpsiLambda_JMM_Filter_DGamma0_TuneCP5_13p6TeV-pythia8-evtgen_cfi',UpgradeFragment(Mby(66,660000),'LbToJpsiLambda_DGamma0_13p6TeV')), #0.3% ('LbToJpsiXiK0sPi_JMM_Filter_DGamma0_TuneCP5_13p6TeV-pythia8-evtgen_cfi',UpgradeFragment(Mby(50,500000),'LbToJpsiXiK0sPr_DGamma0_13p6TeV')), #0.6% ('OmegaMinus_13p6TeV_SoftQCDInel_TuneCP5_cfi',UpgradeFragment(Mby(100,1000000),'OmegaMinus_13p6TeV')), #0.1% + ('Hydjet_Quenched_MinBias_5020GeV_cfi', UpgradeFragment(Kby(9,100),'HydjetQMinBias_5020GeV')), + ('Hydjet_Quenched_MinBias_5362GeV_cfi', UpgradeFragment(Kby(9,100),'HydjetQMinBias_5362GeV')) ]) diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc index a6a26f3fc58fc..2ab9619577039 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc @@ -247,6 +247,9 @@ void SiPixelCompareRecHitsSoA::fillDescriptions(edm::ConfigurationDescription using SiPixelPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA; using SiPixelPhase2CompareRecHitsSoA = SiPixelCompareRecHitsSoA; +using SiPixelHIonPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA; +using SiPixelHIonPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA; DEFINE_FWK_MODULE(SiPixelPhase1CompareRecHitsSoA); DEFINE_FWK_MODULE(SiPixelPhase2CompareRecHitsSoA); +DEFINE_FWK_MODULE(SiPixelHIonPhase1CompareRecHitsSoA); diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc index 3f5e3c6f6bc21..da7ec10f47507 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc @@ -309,6 +309,9 @@ void SiPixelCompareTrackSoA::fillDescriptions(edm::ConfigurationDescriptions& using SiPixelPhase1CompareTrackSoA = SiPixelCompareTrackSoA; using SiPixelPhase2CompareTrackSoA = SiPixelCompareTrackSoA; +using SiPixelHIonPhase1CompareTrackSoA = SiPixelCompareTrackSoA; DEFINE_FWK_MODULE(SiPixelPhase1CompareTrackSoA); DEFINE_FWK_MODULE(SiPixelPhase2CompareTrackSoA); +DEFINE_FWK_MODULE(SiPixelHIonPhase1CompareTrackSoA); + diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorRecHitsSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorRecHitsSoA.cc index 0844bd865ca1f..ac4be07b59ba8 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorRecHitsSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorRecHitsSoA.cc @@ -204,6 +204,9 @@ void SiPixelMonitorRecHitsSoA::fillDescriptions(edm::ConfigurationDescription using SiPixelPhase1MonitorRecHitsSoA = SiPixelMonitorRecHitsSoA; using SiPixelPhase2MonitorRecHitsSoA = SiPixelMonitorRecHitsSoA; +using SiPixelHIonPhase1MonitorRecHitsSoA = SiPixelMonitorRecHitsSoA; DEFINE_FWK_MODULE(SiPixelPhase1MonitorRecHitsSoA); DEFINE_FWK_MODULE(SiPixelPhase2MonitorRecHitsSoA); +DEFINE_FWK_MODULE(SiPixelHIonPhase1MonitorRecHitsSoA); + diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc index 3deb289888477..e971ff184b052 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelMonitorTrackSoA.cc @@ -190,6 +190,8 @@ void SiPixelMonitorTrackSoA::fillDescriptions(edm::ConfigurationDescriptions& using SiPixelPhase1MonitorTrackSoA = SiPixelMonitorTrackSoA; using SiPixelPhase2MonitorTrackSoA = SiPixelMonitorTrackSoA; +using SiPixelHIonPhase1MonitorTrackSoA = SiPixelMonitorTrackSoA; DEFINE_FWK_MODULE(SiPixelPhase1MonitorTrackSoA); DEFINE_FWK_MODULE(SiPixelPhase2MonitorTrackSoA); +DEFINE_FWK_MODULE(SiPixelHIonPhase1MonitorTrackSoA); diff --git a/DQM/SiPixelHeterogeneous/python/SiPixelHeterogenousDQM_FirstStep_cff.py b/DQM/SiPixelHeterogeneous/python/SiPixelHeterogenousDQM_FirstStep_cff.py index 6a263beee2e22..503581e2a2024 100644 --- a/DQM/SiPixelHeterogeneous/python/SiPixelHeterogenousDQM_FirstStep_cff.py +++ b/DQM/SiPixelHeterogeneous/python/SiPixelHeterogenousDQM_FirstStep_cff.py @@ -2,8 +2,10 @@ import FWCore.ParameterSet.Config as cms from DQM.SiPixelHeterogeneous.siPixelPhase1MonitorRecHitsSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase2MonitorRecHitsSoA_cfi import * +from DQM.SiPixelHeterogeneous.siPixelHIonPhase1MonitorRecHitsSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase1MonitorTrackSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase2MonitorTrackSoA_cfi import * +from DQM.SiPixelHeterogeneous.siPixelHIonPhase1MonitorTrackSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelMonitorVertexSoA_cfi import * # Run-3 sequence @@ -14,12 +16,20 @@ _monitorpixelSoARecHitsSource = cms.Sequence(siPixelPhase2MonitorRecHitsSoA * siPixelPhase2MonitorTrackSoA * siPixelMonitorVertexSoA) phase2_tracker.toReplaceWith(monitorpixelSoASource, _monitorpixelSoARecHitsSource) +# HIon Phase 1 sequence +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA + +_monitorpixelSoARecHitsSourceHIon = cms.Sequence(siPixelHIonPhase1MonitorRecHitsSoA * siPixelHIonPhase1MonitorTrackSoA * siPixelMonitorVertexSoA) +pp_on_AA.toReplaceWith(monitorpixelSoASource, _monitorpixelSoARecHitsSourceHIon) + #Define the sequence for GPU vs CPU validation #This should run:- individual monitor for the 2 collections + comparison module from DQM.SiPixelHeterogeneous.siPixelPhase1CompareRecHitsSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase2CompareRecHitsSoA_cfi import * +from DQM.SiPixelHeterogeneous.siPixelHIonPhase1CompareRecHitsSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase1CompareTrackSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase2CompareTrackSoA_cfi import * +from DQM.SiPixelHeterogeneous.siPixelHIonPhase1CompareTrackSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelCompareVertexSoA_cfi import * from DQM.SiPixelHeterogeneous.siPixelPhase1RawDataErrorComparator_cfi import * from DQM.SiPixelPhase1Common.SiPixelPhase1RawData_cfi import * @@ -64,6 +74,16 @@ TopFolderName = "SiPixelHeterogeneous/PixelRecHitsSoAGPU" ) +siPixelHIonPhase1MonitorRecHitsSoACPU = siPixelHIonPhase1MonitorRecHitsSoA.clone( + pixelHitsSrc = "siPixelRecHitsPreSplittingSoA@cpu", + TopFolderName = "SiPixelHeterogeneous/PixelRecHitsSoACPU" +) + +siPixelHIonPhase1MonitorRecHitsSoAGPU = siPixelHIonPhase1MonitorRecHitsSoA.clone( + pixelHitsSrc = "siPixelRecHitsPreSplittingSoA@cuda", + TopFolderName = "SiPixelHeterogeneous/PixelRecHitsSoAGPU" +) + ## tracks siPixelPhase1MonitorTrackSoACPU = siPixelPhase1MonitorTrackSoA.clone( pixelTrackSrc = 'pixelTracksSoA@cpu', @@ -85,6 +105,16 @@ topFolderName = 'SiPixelHeterogeneous/PixelTrackSoAGPU', ) +siPixelHIonPhase1MonitorTrackSoACPU = siPixelHIonPhase1MonitorTrackSoA.clone( + pixelTrackSrc = 'pixelTracksSoA@cpu', + topFolderName = 'SiPixelHeterogeneous/PixelTrackSoACPU', +) + +siPixelHIonPhase1MonitorTrackSoAGPU = siPixelHIonPhase1MonitorTrackSoA.clone( + pixelTrackSrc = 'pixelTracksSoA@cuda', + topFolderName = 'SiPixelHeterogeneous/PixelTrackSoAGPU', +) + ## vertices siPixelMonitorVertexSoACPU = siPixelMonitorVertexSoA.clone( pixelVertexSrc = 'pixelVerticesSoA@cpu', @@ -121,6 +151,17 @@ siPixelMonitorVertexSoAGPU * siPixelCompareVertexSoA) +# HIon sequence +_monitorpixelSoACompareSourceHIonPhase1 = cms.Sequence(siPixelHIonPhase1MonitorRecHitsSoACPU * + siPixelHIonPhase1MonitorRecHitsSoAGPU * + siPixelHIonPhase1CompareRecHitsSoA * + siPixelHIonPhase1MonitorTrackSoAGPU * + siPixelHIonPhase1MonitorTrackSoACPU * + siPixelHIonPhase1CompareTrackSoA * + siPixelMonitorVertexSoACPU * + siPixelMonitorVertexSoAGPU * + siPixelCompareVertexSoA) + phase2_tracker.toReplaceWith(monitorpixelSoACompareSource,_monitorpixelSoACompareSource) from Configuration.ProcessModifiers.gpuValidationPixel_cff import gpuValidationPixel diff --git a/Geometry/CommonTopologies/interface/SimplePixelTopology.h b/Geometry/CommonTopologies/interface/SimplePixelTopology.h index d0e9082122a24..b38a7511882e3 100644 --- a/Geometry/CommonTopologies/interface/SimplePixelTopology.h +++ b/Geometry/CommonTopologies/interface/SimplePixelTopology.h @@ -135,6 +135,8 @@ namespace phase1PixelTopology { constexpr int nPairs = 13 + 2 + 4; constexpr uint16_t numberOfModules = 1856; + constexpr uint32_t maxNumClustersPerModules = 1024; + constexpr uint32_t max_ladder_bpx0 = 12; constexpr uint32_t first_ladder_bpx0 = 0; constexpr float module_length_bpx0 = 6.7f; @@ -208,6 +210,8 @@ namespace phase2PixelTopology { constexpr int nPairs = 23 + 6 + 14 + 8 + 4; // include far forward layer pairs constexpr uint16_t numberOfModules = 3892; + constexpr uint32_t maxNumClustersPerModules = 1024; + HOST_DEVICE_CONSTANT uint8_t layerPairs[2 * nPairs] = { 0, 1, 0, 4, 0, 16, //BPIX1 (3) @@ -282,6 +286,33 @@ namespace phase2PixelTopology { 9.0, 9.0, 8.0, 8.0, 8.0, 11.0, 9.0, 9.0, 9.0, 8.0, 8.0, 8.0, 11.0}; } // namespace phase2PixelTopology +namespace phase1HIonPixelTopology { + using pixelTopology::phi0p09; + + constexpr uint32_t maxNumClustersPerModules = 2048; + + HOST_DEVICE_CONSTANT int16_t phicuts[phase1PixelTopology::nPairs]{phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09, + phi0p09}; + +} // namespace phase1HIonPixelTopology + namespace pixelTopology { struct Phase2 { @@ -317,6 +348,8 @@ namespace pixelTopology { static constexpr uint16_t last_barrel_detIndex = 504; static constexpr uint32_t maxPixInModule = 6000; + static constexpr uint32_t maxNumClustersPerModules = phase2PixelTopology::maxNumClustersPerModules; + static constexpr uint32_t maxHitsInModule = phase2PixelTopology::maxNumClustersPerModules; static constexpr float moduleLength = 4.345f; static constexpr float endcapCorrection = 0.0f; @@ -332,8 +365,6 @@ namespace pixelTopology { static constexpr float bigPixYCorrection = 0.0f; static constexpr float dzdrFact = 8 * 0.0285 / 0.015; // from dz/dr to "DY" - static constexpr float z0Cut = 7.5f; - static constexpr float doubletHardPt = 0.8f; static constexpr int minYsizeB1 = 25; static constexpr int minYsizeB2 = 15; @@ -409,6 +440,8 @@ namespace pixelTopology { static constexpr uint16_t last_barrel_detIndex = 1184; static constexpr uint32_t maxPixInModule = 6000; + static constexpr uint32_t maxNumClustersPerModules = phase1PixelTopology::maxNumClustersPerModules; + static constexpr uint32_t maxHitsInModule = phase1PixelTopology::maxNumClustersPerModules; static constexpr float moduleLength = 6.7f; static constexpr float endcapCorrection = 1.5f; @@ -424,8 +457,6 @@ namespace pixelTopology { static constexpr float bigPixYCorrection = 8.0f; static constexpr float dzdrFact = 8 * 0.0285 / 0.015; // from dz/dr to "DY" - static constexpr float z0Cut = 12.f; - static constexpr float doubletHardPt = 0.5f; static constexpr int minYsizeB1 = 36; static constexpr int minYsizeB2 = 28; @@ -508,15 +539,34 @@ namespace pixelTopology { } }; + struct HIonPhase1 : public Phase1 { + using tindex_type = uint32_t; // for tuples + + static constexpr uint32_t maxCellNeighbors = 90; + static constexpr uint32_t maxCellTracks = 90; + static constexpr uint32_t maxNumberOfTuples = 256 * 1024; + static constexpr uint32_t maxNumberOfDoublets = 6 * 512 * 1024; + static constexpr uint32_t maxHitsForContainers = avgHitsPerTrack * maxNumberOfTuples; + static constexpr uint32_t maxNumberOfQuadruplets = maxNumberOfTuples; + + static constexpr uint32_t maxPixInModule = 10000; + + static constexpr uint32_t maxNumOfActiveDoublets = + maxNumberOfDoublets / 4; //TODO need to think a better way to avoid this duplication + static constexpr uint32_t maxCellsPerHit = 256; + + static constexpr uint32_t maxNumClustersPerModules = phase1HIonPixelTopology::maxNumClustersPerModules; + static constexpr uint32_t maxHitsInModule = phase1HIonPixelTopology::maxNumClustersPerModules; + + static constexpr char const *nameModifier = "HIonPhase1"; + }; + template using isPhase1Topology = typename std::enable_if::value>::type; template using isPhase2Topology = typename std::enable_if::value>::type; - // struct HIonPhase1 : public Phase1 { - // static constexpr uint32_t maxNumberOfDoublets=3*1024*1024;}; - } // namespace pixelTopology #endif // Geometry_CommonTopologies_SimplePixelTopology_h diff --git a/RecoHI/HiTracking/python/HILowPtConformalPixelTracks_cfi.py b/RecoHI/HiTracking/python/HILowPtConformalPixelTracks_cfi.py index 135c927ee75c8..69a3e509b2eed 100644 --- a/RecoHI/HiTracking/python/HILowPtConformalPixelTracks_cfi.py +++ b/RecoHI/HiTracking/python/HILowPtConformalPixelTracks_cfi.py @@ -32,7 +32,7 @@ # Fitter Fitter = 'pixelFitterByConformalMappingAndLine', # Filter - Filter = "hiConformalPixelFilter", + Filter = "hiConformalPixelFilter", # Cleaner Cleaner = "trackCleaner" ) @@ -54,7 +54,7 @@ VertexCollection = "offlinePrimaryVertices", ptMin = 0.3, useFoundVertices = True, - originRadius = 0.2 + originRadius = 0.2 ) ) @@ -62,11 +62,11 @@ # Using 4 layers layerlist from RecoTracker.IterativeTracking.LowPtQuadStep_cff import lowPtQuadStepSeedLayers hiConformalPixelTracksPhase1SeedLayers = lowPtQuadStepSeedLayers.clone( - BPix = cms.PSet( + BPix = cms.PSet( HitProducer = cms.string('siPixelRecHits'), TTRHBuilder = cms.string('WithTrackAngle'), ), - FPix = cms.PSet( + FPix = cms.PSet( HitProducer = cms.string('siPixelRecHits'), TTRHBuilder = cms.string('WithTrackAngle'), ) @@ -86,7 +86,7 @@ doublets = "hiConformalPixelTracksPhase1HitDoubletsCA", CAPhiCut = 0.2, CAThetaCut = 0.0012, - SeedComparitorPSet = dict( + SeedComparitorPSet = dict( ComponentName = 'none' ), extraHitRPhitolerance = 0.032, @@ -130,6 +130,66 @@ hiConformalPixelTracks ) +from Configuration.ProcessModifiers.gpu_cff import gpu +from Configuration.ProcessModifiers.pixelNtupletFit_cff import pixelNtupletFit +from RecoTracker.PixelTrackFitting.pixelTrackSoAFromCUDAHIonPhase1_cfi import pixelTrackSoAFromCUDAHIonPhase1 as _pixelTracksSoA +from RecoTracker.PixelSeeding.caHitNtupletCUDAHIonPhase1_cfi import caHitNtupletCUDAHIonPhase1 as _pixelTracksCUDA +from RecoTracker.PixelTrackFitting.pixelTrackProducerFromSoAHIonPhase1_cfi import pixelTrackProducerFromSoAHIonPhase1 as _pixelTrackProducerFromSoA + +from HeterogeneousCore.CUDACore.SwitchProducerCUDA import SwitchProducerCUDA + +hiPixelTracksCUDA = _pixelTracksCUDA.clone(pixelRecHitSrc="siPixelRecHitsPreSplittingCUDA", idealConditions = False, + ptmin = 0.25, hardCurvCut = 0.0756, doPtCut = False, + onGPU = True, + phiCuts = cms.vint32(19*[900]), #19 pairs + trackQualityCuts = dict( + chi2MaxPt = 10, + chi2Coeff = [0.9,1.8], + chi2Scale = 8, + tripletMinPt = 0.5, + tripletMaxTip = 0.3, + tripletMaxZip = 12, + quadrupletMinPt = 0.3, + quadrupletMaxTip = 0.5, + quadrupletMaxZip = 12 + )) + +# SwitchProducer providing the pixel tracks in SoA format on the CPU +hiPixelTracksSoA = SwitchProducerCUDA( + # build pixel ntuplets and pixel tracks in SoA format on the CPU + cpu = _pixelTracksCUDA.clone( + pixelRecHitSrc = "siPixelRecHitsPreSplittingCPU", + idealConditions = False, + doPtCut = False, + ptmin = 0.25, + hardCurvCut = 0.0756, + onGPU = False, + phiCuts = cms.vint32(19*[900]), #19 pairs + trackQualityCuts = dict( + chi2MaxPt = 10, + chi2Coeff = [0.9,1.8], + chi2Scale = 8, + tripletMinPt = 0.5, + tripletMaxTip = 0.3, + tripletMaxZip = 12, + quadrupletMinPt = 0.3, + quadrupletMaxTip = 0.5, + quadrupletMaxZip = 12 + )) +) + +gpu.toModify(hiPixelTracksSoA, + # transfer the pixel tracks in SoA format to the host + cuda = _pixelTracksSoA.clone(src="hiPixelTracksCUDA") +) + +pixelNtupletFit.toReplaceWith(hiConformalPixelTracks,_pixelTrackProducerFromSoA.clone( + pixelRecHitLegacySrc = "siPixelRecHitsPreSplitting", + trackSrc = "hiPixelTracksSoA", + minQuality = "highPurity" +)) + + hiConformalPixelTracksTaskPhase1 = cms.Task( hiConformalPixelTracksPhase1TrackingRegions , hiConformalPixelTracksPhase1SeedLayers , @@ -139,4 +199,19 @@ hiConformalPixelTracksPhase1Filter , hiConformalPixelTracks ) + +pixelNtupletFit.toReplaceWith(hiConformalPixelTracksTaskPhase1, cms.Task( + # build the pixel ntuplets and the pixel tracks in SoA format on the CPU + hiPixelTracksSoA, + # convert the pixel tracks from SoA to legacy format + hiConformalPixelTracks +)) + +(gpu & pixelNtupletFit).toReplaceWith(hiConformalPixelTracksTaskPhase1, cms.Task( + # build the pixel ntuplets and the pixel tracks in SoA format on the GPU + hiPixelTracksCUDA, + # just copying the task above + hiConformalPixelTracksTaskPhase1.copy() +)) + hiConformalPixelTracksSequencePhase1 = cms.Sequence(hiConformalPixelTracksTaskPhase1) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h index 7cadad85e8329..f9ebb16ea2c7c 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h @@ -5,11 +5,50 @@ struct SiPixelClusterThresholds { inline constexpr int32_t getThresholdForLayerOnCondition(bool isLayer1) const noexcept { return isLayer1 ? layer1 : otherLayers; } - const int32_t layer1; - const int32_t otherLayers; -}; + const int32_t layer1 = 0; + const int32_t otherLayers = 0; + + const float vCaltoElectronGain = 0; + const float vCaltoElectronGain_L1 = 0; + const float vCaltoElectronOffset = 0; + const float vCaltoElectronOffset_L1 = 0; + + const float electronPerADCGain = 0; + const int8_t phase2ReadoutMode = 0; + const uint16_t phase2DigiBaseline = 0; + const uint8_t phase2KinkADC = 0; -constexpr SiPixelClusterThresholds kSiPixelClusterThresholdsDefaultPhase1{.layer1 = 2000, .otherLayers = 4000}; -constexpr SiPixelClusterThresholds kSiPixelClusterThresholdsDefaultPhase2{.layer1 = 4000, .otherLayers = 4000}; + //Basic just for thresholds + SiPixelClusterThresholds(const int32_t layer1, const int32_t otherLayers) + : layer1(layer1), otherLayers(otherLayers) {} + + //For Phase1 + SiPixelClusterThresholds(const int32_t layer1, + const int32_t otherLayers, + const float vCaltoElectronGain, + const float vCaltoElectronGain_L1, + const float vCaltoElectronOffset, + const float vCaltoElectronOffset_L1) + : layer1(layer1), + otherLayers(otherLayers), + vCaltoElectronGain(vCaltoElectronGain), + vCaltoElectronGain_L1(vCaltoElectronGain_L1), + vCaltoElectronOffset(vCaltoElectronOffset), + vCaltoElectronOffset_L1(vCaltoElectronOffset_L1) {} + + //For Phase2 + SiPixelClusterThresholds(const int32_t layer1, + const int32_t otherLayers, + const float electronPerADCGain, + const int8_t phase2ReadoutMode, + const uint16_t phase2DigiBaseline, + const uint8_t phase2KinkADC) + : layer1(layer1), + otherLayers(otherLayers), + electronPerADCGain(electronPerADCGain), + phase2ReadoutMode(phase2ReadoutMode), + phase2DigiBaseline(phase2DigiBaseline), + phase2KinkADC(phase2KinkADC) {} +}; #endif // RecoLocalTracker_SiPixelClusterizer_plugins_SiPixelClusterThresholds_h diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index 3534b6b61d31b..85a5eb161c1fa 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -50,8 +50,8 @@ SiPixelDigisClustersFromSoAT::SiPixelDigisClustersFromSoAT(const : topoToken_(esConsumes()), digiGetToken_(consumes(iConfig.getParameter("src"))), clusterPutToken_(produces()), - clusterThresholds_{iConfig.getParameter("clusterThreshold_layer1"), - iConfig.getParameter("clusterThreshold_otherLayers")}, + clusterThresholds_(iConfig.getParameter("clusterThreshold_layer1"), + iConfig.getParameter("clusterThreshold_otherLayers")), produceDigis_(iConfig.getParameter("produceDigis")), storeDigis_(iConfig.getParameter("produceDigis") && iConfig.getParameter("storeDigis")) { if (produceDigis_) @@ -62,8 +62,8 @@ template void SiPixelDigisClustersFromSoAT::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; desc.add("src", edm::InputTag("siPixelDigisSoA")); - desc.add("clusterThreshold_layer1", kSiPixelClusterThresholdsDefaultPhase1.layer1); - desc.add("clusterThreshold_otherLayers", kSiPixelClusterThresholdsDefaultPhase1.otherLayers); + desc.add("clusterThreshold_layer1", 2000); //TODO put these somewhere + desc.add("clusterThreshold_otherLayers", 4000); desc.add("produceDigis", true); desc.add("storeDigis", true); @@ -108,7 +108,7 @@ void SiPixelDigisClustersFromSoAT::produce(edm::StreamID, } int32_t nclus = -1; - PixelClusterizerBase::AccretionCluster aclusters[gpuClustering::maxNumClustersPerModules]; + PixelClusterizerBase::AccretionCluster aclusters[TrackerTraits::maxNumClustersPerModules]; #ifdef EDM_ML_DEBUG auto totClustersFilled = 0; #endif @@ -183,7 +183,7 @@ void SiPixelDigisClustersFromSoAT::produce(edm::StreamID, // fill clusters #ifdef EDM_ML_DEBUG assert(digis.clus(i) >= 0); - assert(digis.clus(i) < gpuClustering::maxNumClustersPerModules); + assert(digis.clus(i) < TrackerTraits::maxNumClustersPerModules); #endif nclus = std::max(digis.clus(i), nclus); auto row = dig.row(); @@ -209,3 +209,5 @@ using SiPixelDigisClustersFromSoAPhase1 = SiPixelDigisClustersFromSoAT; DEFINE_FWK_MODULE(SiPixelDigisClustersFromSoAPhase2); +using SiPixelDigisClustersFromSoAHIonPhase1 = SiPixelDigisClustersFromSoAT; +DEFINE_FWK_MODULE(SiPixelDigisClustersFromSoAHIonPhase1); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc index 9e19c5ec5ff15..7fac305a9bbd8 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc @@ -33,6 +33,7 @@ #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/ParameterSet/interface/ParameterSetDescription.h" #include "FWCore/ServiceRegistry/interface/Service.h" +#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h" #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "HeterogeneousCore/CUDAServices/interface/CUDAInterface.h" #include "RecoTracker/Record/interface/CkfComponentsRecord.h" @@ -47,6 +48,7 @@ class SiPixelPhase2DigiToClusterCUDA : public edm::stream::EDProducer; private: void acquire(const edm::Event& iEvent, @@ -63,8 +65,7 @@ class SiPixelPhase2DigiToClusterCUDA : public edm::stream::EDProducer wordFedAppender_; + GPUAlgo gpuAlgo_; const bool includeErrors_; const SiPixelClusterThresholds clusterThresholds_; @@ -77,7 +78,11 @@ SiPixelPhase2DigiToClusterCUDA::SiPixelPhase2DigiToClusterCUDA(const edm::Parame clusterPutToken_(produces>()), includeErrors_(iConfig.getParameter("IncludeErrors")), clusterThresholds_{iConfig.getParameter("clusterThreshold_layer1"), - iConfig.getParameter("clusterThreshold_otherLayers")} { + iConfig.getParameter("clusterThreshold_otherLayers"), + (float)iConfig.getParameter("ElectronPerADCGain"), + (int8_t)iConfig.getParameter("Phase2ReadoutMode"), + (uint16_t)iConfig.getParameter("Phase2DigiBaseline"), + (uint8_t)iConfig.getParameter("Phase2KinkADC")} { if (includeErrors_) { digiErrorPutToken_ = produces>(); } @@ -87,8 +92,12 @@ void SiPixelPhase2DigiToClusterCUDA::fillDescriptions(edm::ConfigurationDescript edm::ParameterSetDescription desc; desc.add("IncludeErrors", true); - desc.add("clusterThreshold_layer1", kSiPixelClusterThresholdsDefaultPhase2.layer1); - desc.add("clusterThreshold_otherLayers", kSiPixelClusterThresholdsDefaultPhase2.otherLayers); + desc.add("clusterThreshold_layer1", 4000); + desc.add("clusterThreshold_otherLayers", 4000); + desc.add("ElectronPerADCGain", 1500.f); + desc.add("Phase2ReadoutMode", 3); + desc.add("Phase2DigiBaseline", 1000); + desc.add("Phase2KinkADC", 8); desc.add("InputDigis", edm::InputTag("simSiPixelDigis:Pixel")); descriptions.addWithDefaultLabel(desc); } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc index fd6109958f8ef..aba309f478308 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc @@ -41,12 +41,14 @@ #include "SiPixelClusterThresholds.h" #include "SiPixelRawToClusterGPUKernel.h" -class SiPixelRawToClusterCUDA : public edm::stream::EDProducer { +template +class SiPixelRawToClusterCUDAT : public edm::stream::EDProducer { public: - explicit SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig); - ~SiPixelRawToClusterCUDA() override = default; + explicit SiPixelRawToClusterCUDAT(const edm::ParameterSet& iConfig); + ~SiPixelRawToClusterCUDAT() override = default; static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); + using GPUAlgo = pixelgpudetails::SiPixelRawToClusterGPUKernel; private: void acquire(const edm::Event& iEvent, @@ -72,7 +74,7 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer regions_; - pixelgpudetails::SiPixelRawToClusterGPUKernel gpuAlgo_; + GPUAlgo gpuAlgo_; PixelDataFormatter::Errors errors_; const bool isRun2_; @@ -82,7 +84,8 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer +SiPixelRawToClusterCUDAT::SiPixelRawToClusterCUDAT(const edm::ParameterSet& iConfig) : rawGetToken_(consumes(iConfig.getParameter("InputLabel"))), digiPutToken_(produces>()), clusterPutToken_(produces>()), @@ -94,7 +97,11 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi includeErrors_(iConfig.getParameter("IncludeErrors")), useQuality_(iConfig.getParameter("UseQualityInfo")), clusterThresholds_{iConfig.getParameter("clusterThreshold_layer1"), - iConfig.getParameter("clusterThreshold_otherLayers")} { + iConfig.getParameter("clusterThreshold_otherLayers"), + (float)iConfig.getParameter("VCaltoElectronGain"), + (float)iConfig.getParameter("VCaltoElectronGain_L1"), + (float)iConfig.getParameter("VCaltoElectronOffset"), + (float)iConfig.getParameter("VCaltoElectronOffset_L1")} { if (includeErrors_) { digiErrorPutToken_ = produces>(); } @@ -105,7 +112,8 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi } } -void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { +template +void SiPixelRawToClusterCUDAT::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; desc.add("isRun2", true); desc.add("IncludeErrors", true); @@ -114,8 +122,13 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d // It is kept to avoid breaking older configurations, and will not be printed in the generated cfi.py file. desc.addOptionalNode(edm::ParameterDescription("MaxFEDWords", 0, true), false) ->setComment("This parameter is obsolete and will be ignored."); - desc.add("clusterThreshold_layer1", kSiPixelClusterThresholdsDefaultPhase1.layer1); - desc.add("clusterThreshold_otherLayers", kSiPixelClusterThresholdsDefaultPhase1.otherLayers); + //Clustering Thresholds + desc.add("clusterThreshold_layer1", 2000); + desc.add("clusterThreshold_otherLayers", 4000); + desc.add("VCaltoElectronGain", 47.f); + desc.add("VCaltoElectronGain_L1", 50.f); + desc.add("VCaltoElectronOffset", -60.f); + desc.add("VCaltoElectronOffset_L1", -670.f); desc.add("InputLabel", edm::InputTag("rawDataCollector")); { edm::ParameterSetDescription psd0; @@ -130,9 +143,10 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d descriptions.addWithDefaultLabel(desc); } -void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent, - const edm::EventSetup& iSetup, - edm::WaitingTaskWithArenaHolder waitingTaskHolder) { +template +void SiPixelRawToClusterCUDAT::acquire(const edm::Event& iEvent, + const edm::EventSetup& iSetup, + edm::WaitingTaskWithArenaHolder waitingTaskHolder) { cms::cuda::ScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder), ctxState_}; auto hgpuMap = iSetup.getHandle(gpuMapToken_); @@ -249,27 +263,28 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent, return; // copy the FED data to a single cpu buffer - pixelgpudetails::SiPixelRawToClusterGPUKernel::WordFedAppender wordFedAppender(nDigis_, ctx.stream()); + typename GPUAlgo::WordFedAppender wordFedAppender(nDigis_, ctx.stream()); for (uint32_t i = 0; i < fedIds_.size(); ++i) { wordFedAppender.initializeWordFed(fedIds_[i], index[i], start[i], words[i]); } - gpuAlgo_.makeClustersAsync(isRun2_, - clusterThresholds_, - gpuMap, - gpuModulesToUnpack, - gpuGains, - wordFedAppender, - std::move(errors_), - wordCounter, - fedCounter, - useQuality_, - includeErrors_, - edm::MessageDrop::instance()->debugEnabled, - ctx.stream()); + gpuAlgo_.makePhase1ClustersAsync(isRun2_, + clusterThresholds_, + gpuMap, + gpuModulesToUnpack, + gpuGains, + wordFedAppender, + std::move(errors_), + wordCounter, + fedCounter, + useQuality_, + includeErrors_, + edm::MessageDrop::instance()->debugEnabled, + ctx.stream()); } -void SiPixelRawToClusterCUDA::produce(edm::Event& iEvent, const edm::EventSetup& iSetup) { +template +void SiPixelRawToClusterCUDAT::produce(edm::Event& iEvent, const edm::EventSetup& iSetup) { cms::cuda::ScopedContextProduce ctx{ctxState_}; if (nDigis_ == 0) { @@ -297,4 +312,9 @@ void SiPixelRawToClusterCUDA::produce(edm::Event& iEvent, const edm::EventSetup& } // define as framework plugin +using SiPixelRawToClusterCUDA = SiPixelRawToClusterCUDAT; DEFINE_FWK_MODULE(SiPixelRawToClusterCUDA); +using SiPixelRawToClusterCUDAPhase1 = SiPixelRawToClusterCUDAT; +DEFINE_FWK_MODULE(SiPixelRawToClusterCUDAPhase1); +using SiPixelRawToClusterCUDAHIonPhase1 = SiPixelRawToClusterCUDAT; +DEFINE_FWK_MODULE(SiPixelRawToClusterCUDAHIonPhase1); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 293d4422e8458..419cc1109a732 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -32,6 +32,8 @@ // local includes #include "SiPixelRawToClusterGPUKernel.h" +// #define GPU_DEBUG + namespace pixelgpudetails { __device__ bool isBarrel(uint32_t rawId) { @@ -447,6 +449,8 @@ namespace pixelgpudetails { constexpr int nMaxModules = TrackerTraits::numberOfModules; constexpr int startBPIX2 = TrackerTraits::layerStart[1]; + constexpr uint32_t maxHitsInModule = TrackerTraits::maxHitsInModule; + assert(startBPIX2 < nMaxModules); assert(nMaxModules < 4096); // easy to extend at least till 32*1024 assert(nMaxModules > 1024); @@ -458,7 +462,7 @@ namespace pixelgpudetails { // limit to MaxHitsInModule; for (int i = first, iend = nMaxModules; i < iend; i += blockDim.x) { - moduleStart[i + 1] = std::min(gpuClustering::maxHitsInModule(), clusInModule[i]); + moduleStart[i + 1] = std::min(maxHitsInModule, clusInModule[i]); } constexpr bool isPhase2 = std::is_base_of::value; @@ -500,7 +504,7 @@ namespace pixelgpudetails { #ifdef GPU_DEBUG uint16_t maxH = isPhase2 ? 3027 : 1024; assert(0 == moduleStart[0]); - auto c0 = std::min(gpuClustering::maxHitsInModule(), clusInModule[0]); + auto c0 = std::min(maxHitsInModule, clusInModule[0]); assert(c0 == moduleStart[1]); assert(moduleStart[maxH] >= moduleStart[maxH - 1]); assert(moduleStart[maxH + 1] >= moduleStart[maxH]); @@ -521,20 +525,21 @@ namespace pixelgpudetails { } // Interface to outside - void SiPixelRawToClusterGPUKernel::makeClustersAsync(bool isRun2, - const SiPixelClusterThresholds clusterThresholds, - const SiPixelROCsStatusAndMapping *cablingMap, - const unsigned char *modToUnp, - const SiPixelGainForHLTonGPU *gains, - const WordFedAppender &wordFed, - SiPixelFormatterErrors &&errors, - const uint32_t wordCounter, - const uint32_t fedCounter, - bool useQualityInfo, - bool includeErrors, - bool debug, - cudaStream_t stream) { - using pixelTopology::Phase1; + template + void SiPixelRawToClusterGPUKernel::makePhase1ClustersAsync( + bool isRun2, + const SiPixelClusterThresholds clusterThresholds, + const SiPixelROCsStatusAndMapping *cablingMap, + const unsigned char *modToUnp, + const SiPixelGainForHLTonGPU *gains, + const WordFedAppender &wordFed, + SiPixelFormatterErrors &&errors, + const uint32_t wordCounter, + const uint32_t fedCounter, + bool useQualityInfo, + bool includeErrors, + bool debug, + cudaStream_t stream) { // we're not opting for calling this function in case of early events assert(wordCounter != 0); nDigis = wordCounter; @@ -549,7 +554,7 @@ namespace pixelgpudetails { if (includeErrors) { digiErrors_d = SiPixelDigiErrorsCUDA(wordCounter, std::move(errors), stream); } - clusters_d = SiPixelClustersCUDA(Phase1::numberOfModules, stream); + clusters_d = SiPixelClustersCUDA(TrackerTraits::numberOfModules, stream); // Begin Raw2Digi block { @@ -604,28 +609,19 @@ namespace pixelgpudetails { // clusterizer ... using namespace gpuClustering; int threadsPerBlock = 256; - int blocks = (std::max(int(wordCounter), int(Phase1::numberOfModules)) + threadsPerBlock - 1) / threadsPerBlock; - - if (isRun2) - gpuCalibPixel::calibDigis<<>>(digis_d->moduleId(), - digis_d->xx(), - digis_d->yy(), - digis_d->adc(), - gains, - wordCounter, - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->clusModuleStart()); - else - gpuCalibPixel::calibDigis<<>>(digis_d->moduleId(), - digis_d->xx(), - digis_d->yy(), - digis_d->adc(), - gains, - wordCounter, - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->clusModuleStart()); + int blocks = + (std::max(int(wordCounter), int(TrackerTraits::numberOfModules)) + threadsPerBlock - 1) / threadsPerBlock; + + gpuCalibPixel::calibDigis<<>>(clusterThresholds, + digis_d.view().moduleId(), + digis_d.view().xx(), + digis_d.view().yy(), + digis_d.view().adc(), + gains, + wordCounter, + clusters_d->moduleStart(), + clusters_d->clusInModule(), + clusters_d->clusModuleStart()); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG @@ -637,25 +633,26 @@ namespace pixelgpudetails { << " threads\n"; #endif - countModules<<>>( + countModules<<>>( digis_d->moduleId(), clusters_d->moduleStart(), digis_d->clus(), wordCounter); cudaCheck(cudaGetLastError()); - threadsPerBlock = 256 + 128; /// should be larger than 6000/16 aka (maxPixInModule/maxiter in the kernel) - blocks = phase1PixelTopology::numberOfModules; + threadsPerBlock = ((TrackerTraits::maxPixInModule / 16 + 128 - 1) / 128) * + 128; /// should be larger than maxPixInModule/16 aka (maxPixInModule/maxiter in the kernel) + blocks = TrackerTraits::numberOfModules; #ifdef GPU_DEBUG std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif - findClus<<>>(digis_d->rawIdArr(), - digis_d->moduleId(), - digis_d->xx(), - digis_d->yy(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), - wordCounter); + findClus<<>>(digis_d->rawIdArr(), + digis_d->moduleId(), + digis_d->xx(), + digis_d->yy(), + clusters_d->moduleStart(), + clusters_d->clusInModule(), + clusters_d->moduleId(), + digis_d->clus(), + wordCounter); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG @@ -663,14 +660,14 @@ namespace pixelgpudetails { #endif // apply charge cut - clusterChargeCut<<>>(clusterThresholds, - digis_d->moduleId(), - digis_d->adc(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), - wordCounter); + clusterChargeCut<<>>(clusterThresholds, + digis_d->moduleId(), + digis_d->adc(), + clusters_d->moduleStart(), + clusters_d->clusInModule(), + clusters_d->moduleId(), + digis_d->clus(), + wordCounter); cudaCheck(cudaGetLastError()); @@ -680,10 +677,10 @@ namespace pixelgpudetails { // synchronization/ExternalWork auto nModules_Clusters_d = cms::cuda::make_device_unique(3, stream); // MUST be ONE block - fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule(), - clusters_d->clusModuleStart(), - clusters_d->moduleStart(), - nModules_Clusters_d.get()); + fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule(), + clusters_d->clusModuleStart(), + clusters_d->moduleStart(), + nModules_Clusters_d.get()); // copy to host nModules_Clusters_h = cms::cuda::make_host_unique(3, stream); @@ -697,17 +694,18 @@ namespace pixelgpudetails { } // end clusterizer scope } - void SiPixelRawToClusterGPUKernel::makePhase2ClustersAsync(const SiPixelClusterThresholds clusterThresholds, - const uint16_t *moduleIds, - const uint16_t *xDigis, - const uint16_t *yDigis, - const uint16_t *adcDigis, - const uint32_t *packedData, - const uint32_t *rawIds, - const uint32_t numDigis, - cudaStream_t stream) { + template + void SiPixelRawToClusterGPUKernel::makePhase2ClustersAsync( + const SiPixelClusterThresholds clusterThresholds, + const uint16_t *moduleIds, + const uint16_t *xDigis, + const uint16_t *yDigis, + const uint16_t *adcDigis, + const uint32_t *packedData, + const uint32_t *rawIds, + const uint32_t numDigis, + cudaStream_t stream) { using namespace gpuClustering; - using pixelTopology::Phase2; nDigis = numDigis; digis_d = SiPixelDigisCUDA(numDigis, stream); @@ -718,14 +716,15 @@ namespace pixelgpudetails { cudaCheck(cudaMemcpyAsync(digis_d->pdigi(), packedData, sizeof(uint32_t) * numDigis, cudaMemcpyDefault, stream)); cudaCheck(cudaMemcpyAsync(digis_d->rawIdArr(), rawIds, sizeof(uint32_t) * numDigis, cudaMemcpyDefault, stream)); - clusters_d = SiPixelClustersCUDA(Phase2::numberOfModules, stream); + clusters_d = SiPixelClustersCUDA(TrackerTraits::numberOfModules, stream); nModules_Clusters_h = cms::cuda::make_host_unique(2, stream); int threadsPerBlock = 512; int blocks = (int(numDigis) + threadsPerBlock - 1) / threadsPerBlock; - gpuCalibPixel::calibDigisPhase2<<>>(digis_d->moduleId(), + gpuCalibPixel::calibDigisPhase2<<>>(clusterThresholds, + digis_d->moduleId(), digis_d->adc(), numDigis, clusters_d->moduleStart(), @@ -739,7 +738,7 @@ namespace pixelgpudetails { std::cout << "CUDA countModules kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif - countModules<<>>( + countModules<<>>( digis_d->moduleId(), clusters_d->moduleStart(), digis_d->clus(), numDigis); cudaCheck(cudaGetLastError()); @@ -748,21 +747,21 @@ namespace pixelgpudetails { &(nModules_Clusters_h[0]), clusters_d->moduleStart(), sizeof(uint32_t), cudaMemcpyDefault, stream)); threadsPerBlock = 256; - blocks = Phase2::numberOfModules; + blocks = TrackerTraits::numberOfModules; #ifdef GPU_DEBUG cudaCheck(cudaStreamSynchronize(stream)); std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif - findClus<<>>(digis_d->rawIdArr(), - digis_d->moduleId(), - digis_d->xx(), - digis_d->yy(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), - numDigis); + findClus<<>>(digis_d->rawIdArr(), + digis_d->moduleId(), + digis_d->xx(), + digis_d->yy(), + clusters_d->moduleStart(), + clusters_d->clusInModule(), + clusters_d->moduleId(), + digis_d->clus(), + numDigis); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG @@ -772,14 +771,14 @@ namespace pixelgpudetails { #endif // apply charge cut - clusterChargeCut<<>>(clusterThresholds, - digis_d->moduleId(), - digis_d->adc(), - clusters_d->moduleStart(), - clusters_d->clusInModule(), - clusters_d->moduleId(), - digis_d->clus(), - numDigis); + clusterChargeCut<<>>(clusterThresholds, + digis_d->moduleId(), + digis_d->adc(), + clusters_d->moduleStart(), + clusters_d->clusInModule(), + clusters_d->moduleId(), + digis_d->clus(), + numDigis); cudaCheck(cudaGetLastError()); auto nModules_Clusters_d = cms::cuda::make_device_unique(3, stream); @@ -790,10 +789,10 @@ namespace pixelgpudetails { std::cout << "CUDA fillHitsModuleStart kernel launch \n"; #endif - fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule(), - clusters_d->clusModuleStart(), - clusters_d->moduleStart(), - nModules_Clusters_d.get()); + fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d->clusInModule(), + clusters_d->clusModuleStart(), + clusters_d->moduleStart(), + nModules_Clusters_d.get()); nModules_Clusters_h = cms::cuda::make_host_unique(3, stream); cudaCheck(cudaMemcpyAsync( @@ -803,4 +802,8 @@ namespace pixelgpudetails { cudaCheck(cudaStreamSynchronize(stream)); #endif } // + + template class SiPixelRawToClusterGPUKernel; + template class SiPixelRawToClusterGPUKernel; + template class SiPixelRawToClusterGPUKernel; } // namespace pixelgpudetails diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index ace787514486e..98d079816ddd3 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h @@ -5,6 +5,8 @@ #include #include "DataFormats/SiPixelDetId/interface/PixelChannelIdentifier.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h" #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h" @@ -12,8 +14,9 @@ #include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" -#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h" -#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h" +#include "Geometry/CommonTopologies/interface/SimplePixelTopology.h" + +// #define GPU_DEBUG // local include(s) #include "SiPixelClusterThresholds.h" @@ -71,6 +74,7 @@ namespace pixelgpudetails { return (row << thePacking.column_width) | col; } + template class SiPixelRawToClusterGPUKernel { public: class WordFedAppender { @@ -100,19 +104,19 @@ namespace pixelgpudetails { SiPixelRawToClusterGPUKernel& operator=(const SiPixelRawToClusterGPUKernel&) = delete; SiPixelRawToClusterGPUKernel& operator=(SiPixelRawToClusterGPUKernel&&) = delete; - void makeClustersAsync(bool isRun2, - const SiPixelClusterThresholds clusterThresholds, - const SiPixelROCsStatusAndMapping* cablingMap, - const unsigned char* modToUnp, - const SiPixelGainForHLTonGPU* gains, - const WordFedAppender& wordFed, - SiPixelFormatterErrors&& errors, - const uint32_t wordCounter, - const uint32_t fedCounter, - bool useQualityInfo, - bool includeErrors, - bool debug, - cudaStream_t stream); + void makePhase1ClustersAsync(bool isRun2, + const SiPixelClusterThresholds clusterThresholds, + const SiPixelROCsStatusAndMapping* cablingMap, + const unsigned char* modToUnp, + const SiPixelGainForHLTonGPU* gains, + const WordFedAppender& wordFed, + SiPixelFormatterErrors&& errors, + const uint32_t wordCounter, + const uint32_t fedCounter, + bool useQualityInfo, + bool includeErrors, + bool debug, + cudaStream_t stream); void makePhase2ClustersAsync(const SiPixelClusterThresholds clusterThresholds, const uint16_t* moduleIds, diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h index 75e8389513b68..180b356db2c88 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h @@ -11,25 +11,16 @@ #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "Geometry/CommonTopologies/interface/SimplePixelTopology.h" +// local include(s) +#include "SiPixelClusterThresholds.h" + namespace gpuCalibPixel { using gpuClustering::invalidModuleId; - // calibrationConstants - // valid for run2 - constexpr float VCaltoElectronGain = 47; // L2-4: 47 +- 4.7 - constexpr float VCaltoElectronGain_L1 = 50; // L1: 49.6 +- 2.6 - constexpr float VCaltoElectronOffset = -60; // L2-4: -60 +- 130 - constexpr float VCaltoElectronOffset_L1 = -670; // L1: -670 +- 220 - constexpr int VCalChargeThreshold = 100; - //for phase2 - constexpr float ElectronPerADCGain = 1500; - constexpr int8_t Phase2ReadoutMode = 3; - constexpr uint16_t Phase2DigiBaseline = 1000; - constexpr uint8_t Phase2KinkADC = 8; - - template - __global__ void calibDigis(uint16_t* id, + // template + __global__ void calibDigis(SiPixelClusterThresholds clusterThresholds, + uint16_t* id, uint16_t const* __restrict__ x, uint16_t const* __restrict__ y, uint16_t* adc, @@ -41,6 +32,11 @@ namespace gpuCalibPixel { ) { int first = blockDim.x * blockIdx.x + threadIdx.x; + const float VCaltoElectronGain = clusterThresholds.vCaltoElectronGain; + const float VCaltoElectronGain_L1 = clusterThresholds.vCaltoElectronGain_L1; + const float VCaltoElectronOffset = clusterThresholds.vCaltoElectronOffset; + const float VCaltoElectronOffset_L1 = clusterThresholds.vCaltoElectronOffset_L1; + // zero for next kernels... if (0 == first) clusModuleStart[0] = moduleStart[0] = 0; @@ -66,17 +62,18 @@ namespace gpuCalibPixel { adc[i] = 0; } else { float vcal = float(adc[i]) * gain - pedestal * gain; - if constexpr (isRun2) { - float conversionFactor = id[i] < 96 ? VCaltoElectronGain_L1 : VCaltoElectronGain; - float offset = id[i] < 96 ? VCaltoElectronOffset_L1 : VCaltoElectronOffset; - vcal = vcal * conversionFactor + offset; - } + + float conversionFactor = id[i] < 96 ? VCaltoElectronGain_L1 : VCaltoElectronGain; + float offset = id[i] < 96 ? VCaltoElectronOffset_L1 : VCaltoElectronOffset; + vcal = vcal * conversionFactor + offset; + adc[i] = std::clamp(int(vcal), 100, int(std::numeric_limits::max())); } } } - __global__ void calibDigisPhase2(uint16_t* id, + __global__ void calibDigisPhase2(SiPixelClusterThresholds clusterThresholds, + uint16_t* id, uint16_t* adc, int numElements, uint32_t* __restrict__ moduleStart, // just to zero first @@ -86,6 +83,11 @@ namespace gpuCalibPixel { int first = blockDim.x * blockIdx.x + threadIdx.x; // zero for next kernels... + const float ElectronPerADCGain = clusterThresholds.electronPerADCGain; + const int8_t Phase2ReadoutMode = clusterThresholds.phase2ReadoutMode; + const uint16_t Phase2DigiBaseline = clusterThresholds.phase2DigiBaseline; + const uint8_t Phase2KinkADC = clusterThresholds.phase2KinkADC; + if (0 == first) clusModuleStart[0] = moduleStart[0] = 0; for (int i = first; i < phase2PixelTopology::numberOfModules; i += gridDim.x * blockDim.x) { @@ -96,18 +98,18 @@ namespace gpuCalibPixel { if (invalidModuleId == id[i]) continue; - constexpr int mode = (Phase2ReadoutMode < -1 ? -1 : Phase2ReadoutMode); + const int mode = (Phase2ReadoutMode < -1 ? -1 : Phase2ReadoutMode); int adc_int = adc[i]; - if constexpr (mode < 0) + if (mode < 0) adc_int = int(adc_int * ElectronPerADCGain); else { if (adc_int < Phase2KinkADC) adc_int = int((adc_int + 0.5) * ElectronPerADCGain); else { - constexpr int8_t dspp = (Phase2ReadoutMode < 10 ? Phase2ReadoutMode : 10); - constexpr int8_t ds = int8_t(dspp <= 1 ? 1 : (dspp - 1) * (dspp - 1)); + const int8_t dspp = (Phase2ReadoutMode < 10 ? Phase2ReadoutMode : 10); + const int8_t ds = int8_t(dspp <= 1 ? 1 : (dspp - 1) * (dspp - 1)); adc_int -= Phase2KinkADC; adc_int *= ds; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h index fced5675e5c29..cfd6efb3eef2d 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -25,6 +25,13 @@ namespace gpuClustering { uint32_t const* __restrict__ moduleId, // module id of each module int32_t* __restrict__ clusterId, // modified: cluster id of each pixel uint32_t numElements) { + constexpr int32_t maxNumClustersPerModules = TrackerTraits::maxNumClustersPerModules; + + static_assert( + maxNumClustersPerModules <= 2048, + "\nclusterChargeCut is limited to 2048 clusters per module. \nHere maxNumClustersPerModules is set to be %d. " + "\nIf you need maxNumClustersPerModules to be higher \nyou will need to fix the blockPrefixScans."); + __shared__ int32_t charge[maxNumClustersPerModules]; __shared__ uint8_t ok[maxNumClustersPerModules]; __shared__ uint16_t newclusId[maxNumClustersPerModules]; @@ -119,8 +126,21 @@ namespace gpuClustering { // renumber __shared__ uint16_t ws[32]; - cms::cuda::blockPrefixScan(newclusId, nclus, ws); - + constexpr auto maxThreads = 1024; + auto minClust = nclus > maxThreads ? maxThreads : nclus; + + cms::cuda::blockPrefixScan(newclusId, newclusId, minClust, ws); + if constexpr (maxNumClustersPerModules > maxThreads) //only if needed + { + //TODO: most probably there's a smarter implementation for this + if (nclus > maxThreads) { + cms::cuda::blockPrefixScan(newclusId + maxThreads, newclusId + maxThreads, nclus - maxThreads, ws); + for (auto i = threadIdx.x + maxThreads; i < nclus; i += blockDim.x) { + int prevBlockEnd = ((i / maxThreads) * maxThreads) - 1; + newclusId[i] += newclusId[prevBlockEnd]; + } + } + } assert(nclus > newclusId[nclus - 1]); nClustersInModule[thisModuleId] = newclusId[nclus - 1]; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h index 675eae8938236..1a9395b8e7229 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h @@ -163,6 +163,9 @@ namespace gpuClustering { printf("too many pixels in module %d: %d > %d\n", thisModuleId, msize - firstPixel, maxPixInModule); msize = maxPixInModule + firstPixel; } +#ifdef GPU_DEBUG + printf("pixelInModule > %d\n", msize - firstPixel); +#endif } __syncthreads(); diff --git a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py index 66b924b72d1c2..2f193f7c861f4 100644 --- a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py +++ b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py @@ -15,20 +15,34 @@ ) # reconstruct the pixel digis and clusters on the gpu -from RecoLocalTracker.SiPixelClusterizer.siPixelRawToClusterCUDA_cfi import siPixelRawToClusterCUDA as _siPixelRawToClusterCUDA +from RecoLocalTracker.SiPixelClusterizer.siPixelRawToClusterCUDAPhase1_cfi import siPixelRawToClusterCUDAPhase1 as _siPixelRawToClusterCUDA +from RecoLocalTracker.SiPixelClusterizer.siPixelRawToClusterCUDAHIonPhase1_cfi import siPixelRawToClusterCUDAHIonPhase1 as _siPixelRawToClusterCUDAHIonPhase1 siPixelClustersPreSplittingCUDA = _siPixelRawToClusterCUDA.clone() +# HIon modifiers +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA + +pp_on_AA.toReplaceWith(siPixelClustersPreSplittingCUDA, _siPixelRawToClusterCUDAHIonPhase1.clone()) + run3_common.toModify(siPixelClustersPreSplittingCUDA, # use the pixel channel calibrations scheme for Run 3 isRun2 = False, - clusterThreshold_layer1 = 4000) + clusterThreshold_layer1 = 4000, + VCaltoElectronGain = 1, # all gains=1, pedestals=0 + VCaltoElectronGain_L1 = 1, + VCaltoElectronOffset = 0, + VCaltoElectronOffset_L1 = 0) + -# convert the pixel digis (except errors) and clusters to the legacy format from RecoLocalTracker.SiPixelClusterizer.siPixelDigisClustersFromSoAPhase1_cfi import siPixelDigisClustersFromSoAPhase1 as _siPixelDigisClustersFromSoAPhase1 from RecoLocalTracker.SiPixelClusterizer.siPixelDigisClustersFromSoAPhase2_cfi import siPixelDigisClustersFromSoAPhase2 as _siPixelDigisClustersFromSoAPhase2 siPixelDigisClustersPreSplitting = _siPixelDigisClustersFromSoAPhase1.clone() +from RecoLocalTracker.SiPixelClusterizer.siPixelDigisClustersFromSoAHIonPhase1_cfi import siPixelDigisClustersFromSoAHIonPhase1 as _siPixelDigisClustersFromSoAHIonPhase1 +pp_on_AA.toReplaceWith(siPixelDigisClustersPreSplitting, _siPixelDigisClustersFromSoAHIonPhase1.clone()) + + run3_common.toModify(siPixelDigisClustersPreSplitting, clusterThreshold_layer1 = 4000) @@ -50,7 +64,14 @@ # for phase2 no pixel raw2digi is available at the moment # so we skip the raw2digi step and run on pixel digis copied to gpu -phase2_tracker.toReplaceWith(siPixelClustersPreSplittingCUDA,_siPixelPhase2DigiToClusterCUDA.clone()) +from SimTracker.SiPhase2Digitizer.phase2TrackerDigitizer_cfi import PixelDigitizerAlgorithmCommon + +phase2_tracker.toReplaceWith(siPixelClustersPreSplittingCUDA,_siPixelPhase2DigiToClusterCUDA.clone( + Phase2ReadoutMode = PixelDigitizerAlgorithmCommon.Phase2ReadoutMode.value(), # Flag to decide Readout Mode : linear TDR (-1), dual slope with slope parameters (+1,+2,+3,+4 ...) with threshold subtraction + Phase2DigiBaseline = int(PixelDigitizerAlgorithmCommon.ThresholdInElectrons_Barrel.value()), #Same for barrel and endcap + Phase2KinkADC = 8, + ElectronPerADCGain = PixelDigitizerAlgorithmCommon.ElectronPerAdc.value() +)) from EventFilter.SiPixelRawToDigi.siPixelDigisSoAFromCUDA_cfi import siPixelDigisSoAFromCUDA as _siPixelDigisSoAFromCUDA siPixelDigisPhase2SoA = _siPixelDigisSoAFromCUDA.clone( diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h index 6aff7aa15196e..a8f7dae43b5d7 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h +++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h @@ -31,7 +31,7 @@ int main(void) { using pixelTopology::Phase1; constexpr int numElements = 256 * maxNumModules; - constexpr SiPixelClusterThresholds clusterThresholds(kSiPixelClusterThresholdsDefaultPhase1); + const SiPixelClusterThresholds clusterThresholds(2000, 4000, 0.f, 0.f, 0.f, 0.f); // these in reality are already on GPU auto h_raw = std::make_unique(numElements); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc b/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc index 3e9a128f7315c..171cfd1baad55 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc @@ -106,3 +106,5 @@ using PixelCPEFastESProducerPhase1 = PixelCPEFastESProducerT; DEFINE_FWK_EVENTSETUP_MODULE(PixelCPEFastESProducerPhase2); +using PixelCPEFastESProducerHIonPhase1 = PixelCPEFastESProducerT; +DEFINE_FWK_EVENTSETUP_MODULE(PixelCPEFastESProducerHIonPhase1); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu index e2d2b06344308..61442ea9d2b8c 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu @@ -102,4 +102,5 @@ namespace pixelgpudetails { template class PixelRecHitGPUKernel; template class PixelRecHitGPUKernel; + template class PixelRecHitGPUKernel; } // namespace pixelgpudetails diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc index cbed0bfa22cdb..6a5364beed69a 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc @@ -99,3 +99,6 @@ DEFINE_FWK_MODULE(SiPixelRecHitCUDAPhase1); using SiPixelRecHitCUDAPhase2 = SiPixelRecHitCUDAT; DEFINE_FWK_MODULE(SiPixelRecHitCUDAPhase2); + +using SiPixelRecHitCUDAHIonPhase1 = SiPixelRecHitCUDAT; +DEFINE_FWK_MODULE(SiPixelRecHitCUDAHIonPhase1); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc index ba42c75200623..8ef3f74da5751 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc @@ -119,7 +119,7 @@ void SiPixelRecHitFromCUDAT::produce(edm::Event& iEvent, edm::Eve edm::Handle hclusters = iEvent.getHandle(clusterToken_); auto const& input = *hclusters; - constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule(); + constexpr uint32_t maxHitsInModule = TrackerTraits::maxHitsInModule; int numberOfDetUnits = 0; int numberOfClusters = 0; @@ -198,3 +198,6 @@ DEFINE_FWK_MODULE(SiPixelRecHitFromCUDAPhase1); using SiPixelRecHitFromCUDAPhase2 = SiPixelRecHitFromCUDAT; DEFINE_FWK_MODULE(SiPixelRecHitFromCUDAPhase2); + +using SiPixelRecHitFromCUDAHIonPhase1 = SiPixelRecHitFromCUDAT; +DEFINE_FWK_MODULE(SiPixelRecHitFromCUDAHIonPhase1); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromCUDA.cc index df42c753c5632..c9ba2728243a6 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromCUDA.cc @@ -98,3 +98,6 @@ DEFINE_FWK_MODULE(SiPixelRecHitSoAFromCUDAPhase1); using SiPixelRecHitSoAFromCUDAPhase2 = SiPixelRecHitSoAFromCUDAT; DEFINE_FWK_MODULE(SiPixelRecHitSoAFromCUDAPhase2); + +using SiPixelRecHitSoAFromCUDAHIonPhase1 = SiPixelRecHitSoAFromCUDAT; +DEFINE_FWK_MODULE(SiPixelRecHitSoAFromCUDAHIonPhase1); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 678119420edbb..1da748d8dcd8b 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -291,3 +291,6 @@ DEFINE_FWK_MODULE(SiPixelRecHitSoAFromLegacyPhase1); using SiPixelRecHitSoAFromLegacyPhase2 = SiPixelRecHitSoAFromLegacyT; DEFINE_FWK_MODULE(SiPixelRecHitSoAFromLegacyPhase2); + +using SiPixelRecHitSoAFromLegacyHIonPhase1 = SiPixelRecHitSoAFromLegacyT; +DEFINE_FWK_MODULE(SiPixelRecHitSoAFromLegacyHIonPhase1); diff --git a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py index 4c0e007847a03..686b0afc335c4 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py +++ b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py @@ -12,6 +12,7 @@ from RecoLocalTracker.SiPixelRecHits.PixelCPEGeneric_cfi import * from RecoLocalTracker.SiPixelRecHits.pixelCPEFastESProducerPhase1_cfi import * from RecoLocalTracker.SiPixelRecHits.pixelCPEFastESProducerPhase2_cfi import * +from RecoLocalTracker.SiPixelRecHits.pixelCPEFastESProducerHIonPhase1_cfi import * # # 3. ESProducer for the Magnetic-field dependent template records # diff --git a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py index 68f3869c9b458..c4ee471950b93 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py +++ b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py @@ -18,12 +18,16 @@ # phase 2 tracker modifier from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker +# HIon modifiers +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA # convert the pixel rechits from legacy to SoA format on CPU from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromLegacyPhase1_cfi import siPixelRecHitSoAFromLegacyPhase1 as _siPixelRecHitsPreSplittingSoA from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromLegacyPhase2_cfi import siPixelRecHitSoAFromLegacyPhase2 as _siPixelRecHitsPreSplittingSoAPhase2 +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromLegacyHIonPhase1_cfi import siPixelRecHitSoAFromLegacyHIonPhase1 as _siPixelRecHitsPreSplittingSoAHIonPhase1 siPixelRecHitsPreSplittingCPU = _siPixelRecHitsPreSplittingSoA.clone(convertToLegacy=True) +pp_on_AA.toReplaceWith(siPixelRecHitsPreSplittingCPU, _siPixelRecHitsPreSplittingSoAHIonPhase1.clone(convertToLegacy=True, CPE = cms.string('PixelCPEFastHIonPhase1'))) phase2_tracker.toReplaceWith(siPixelRecHitsPreSplittingCPU, _siPixelRecHitsPreSplittingSoAPhase2.clone(convertToLegacy=True, CPE = cms.string('PixelCPEFastPhase2'))) # modifier used to prompt patatrack pixel tracks reconstruction on cpu @@ -43,11 +47,17 @@ ) # reconstruct the pixel rechits on the gpu + from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDAPhase1_cfi import siPixelRecHitCUDAPhase1 as _siPixelRecHitCUDAPhase1 from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDAPhase2_cfi import siPixelRecHitCUDAPhase2 as _siPixelRecHitCUDAPhase2 +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDAHIonPhase1_cfi import siPixelRecHitCUDAHIonPhase1 as _siPixelRecHitCUDAHIonPhase1 + siPixelRecHitsPreSplittingCUDA = _siPixelRecHitCUDAPhase1.clone( beamSpot = "offlineBeamSpotToCUDA" ) +pp_on_AA.toReplaceWith(siPixelRecHitsPreSplittingCUDA,_siPixelRecHitCUDAHIonPhase1.clone( + beamSpot = "offlineBeamSpotToCUDA" +)) phase2_tracker.toReplaceWith(siPixelRecHitsPreSplittingCUDA,_siPixelRecHitCUDAPhase2.clone( beamSpot = "offlineBeamSpotToCUDA" )) @@ -61,6 +71,13 @@ )), ) +pp_on_AA.toModify(siPixelRecHitsPreSplittingSoA, +cpu = cms.EDAlias( + siPixelRecHitsPreSplittingCPU = cms.VPSet( + cms.PSet(type = cms.string("pixelTopologyHIonPhase1TrackingRecHitSoAHost")), + cms.PSet(type = cms.string("uintAsHostProduct")) + ))) + phase2_tracker.toModify(siPixelRecHitsPreSplittingSoA, cpu = cms.EDAlias( siPixelRecHitsPreSplittingCPU = cms.VPSet( @@ -69,6 +86,7 @@ ))) from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromCUDAPhase1_cfi import siPixelRecHitSoAFromCUDAPhase1 as _siPixelRecHitSoAFromCUDA +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromCUDAHIonPhase1_cfi import siPixelRecHitSoAFromCUDAHIonPhase1 as _siPixelRecHitSoAFromCUDAHIonPhase1 from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromCUDAPhase2_cfi import siPixelRecHitSoAFromCUDAPhase2 as _siPixelRecHitSoAFromCUDAPhase2 (gpu & pixelNtupletFit).toModify(siPixelRecHitsPreSplittingSoA, cuda = _siPixelRecHitSoAFromCUDA.clone()) @@ -76,13 +94,14 @@ # transfer the pixel rechits to the host and convert them from SoA from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromCUDAPhase1_cfi import siPixelRecHitFromCUDAPhase1 as _siPixelRecHitFromCUDA +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromCUDAHIonPhase1_cfi import siPixelRecHitFromCUDAHIonPhase1 as _siPixelRecHitFromCUDAHIonPhase1 from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromCUDAPhase2_cfi import siPixelRecHitFromCUDAPhase2 as _siPixelRecHitFromCUDAPhase2 (gpu & pixelNtupletFit).toModify(siPixelRecHitsPreSplitting, cuda = _siPixelRecHitFromCUDA.clone()) +(gpu & pixelNtupletFit & pp_on_AA).toModify(siPixelRecHitsPreSplitting, cuda = _siPixelRecHitFromCUDAHIonPhase1.clone()) (gpu & pixelNtupletFit & phase2_tracker).toModify(siPixelRecHitsPreSplitting, cuda = _siPixelRecHitFromCUDAPhase2.clone()) - pixelNtupletFit.toReplaceWith(siPixelRecHitsPreSplittingTask, cms.Task( cms.Task( # reconstruct the pixel rechits on the cpu diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc index 77f8b56f47c1d..e2e63b56a35df 100644 --- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc +++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc @@ -536,3 +536,4 @@ void PixelCPEFast::fillPSetDescription(edm::ParameterSetDescripti template class PixelCPEFast; template class PixelCPEFast; +template class PixelCPEFast; diff --git a/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py b/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py index 93a5cfa401ad3..0a3e57a9ff96a 100644 --- a/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py +++ b/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py @@ -13,9 +13,13 @@ from Configuration.ProcessModifiers.pixelNtupletFit_cff import pixelNtupletFit from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker +# HIon modifiers +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA + # build the pixel vertices in SoA format on the CPU from RecoTracker.PixelVertexFinding.pixelVertexProducerCUDAPhase1_cfi import pixelVertexProducerCUDAPhase1 as _pixelVerticesCUDA from RecoTracker.PixelVertexFinding.pixelVertexProducerCUDAPhase2_cfi import pixelVertexProducerCUDAPhase2 as _pixelVerticesCUDAPhase2 +from RecoTracker.PixelVertexFinding.pixelVertexProducerCUDAHIonPhase1_cfi import pixelVertexProducerCUDAHIonPhase1 as _pixelVerticesCUDAHIonPhase1 pixelVerticesSoA = SwitchProducerCUDA( cpu = _pixelVerticesCUDA.clone( @@ -30,6 +34,12 @@ PtMin = 2.0 )) +pp_on_AA.toModify(pixelVerticesSoA,cpu = _pixelVerticesCUDAHIonPhase1.clone( + pixelTrackSrc = "pixelTracksSoA", + doSplitting = False, + onGPU = False, +)) + # convert the pixel vertices from SoA to legacy format from RecoTracker.PixelVertexFinding.pixelVertexFromSoA_cfi import pixelVertexFromSoA as _pixelVertexFromSoA @@ -60,6 +70,12 @@ PtMin = 2.0 )) +pp_on_AA.toReplaceWith(pixelVerticesCUDA,_pixelVerticesCUDAHIonPhase1.clone( + pixelTrackSrc = "pixelTracksCUDA", + doSplitting = False, + onGPU = True +)) + # transfer the pixel vertices in SoA format to the CPU from RecoTracker.PixelVertexFinding.pixelVerticesSoA_cfi import pixelVerticesSoA as _pixelVerticesSoA gpu.toModify(pixelVerticesSoA, diff --git a/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cc b/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cc index be92f2d5d0fa2..ad3de7be225db 100644 --- a/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cc +++ b/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cc @@ -115,3 +115,4 @@ void HelixFitOnGPU::launchBrokenLineKernelsOnCPU(const TrackingRe template class HelixFitOnGPU; template class HelixFitOnGPU; +template class HelixFitOnGPU; diff --git a/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cu b/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cu index c5c9ac7fc6345..1558fa9ae6176 100644 --- a/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cu +++ b/RecoTracker/PixelSeeding/plugins/BrokenLineFitOnGPU.cu @@ -142,3 +142,4 @@ void HelixFitOnGPU::launchBrokenLineKernels(const TrackingRecHitS template class HelixFitOnGPU; template class HelixFitOnGPU; +template class HelixFitOnGPU; diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletCUDA.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletCUDA.cc index d63fb59fec8e0..06ca6476bc4fa 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletCUDA.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletCUDA.cc @@ -113,3 +113,6 @@ DEFINE_FWK_MODULE(CAHitNtupletCUDAPhase1); using CAHitNtupletCUDAPhase2 = CAHitNtupletCUDAT; DEFINE_FWK_MODULE(CAHitNtupletCUDAPhase2); + +using CAHitNtupletCUDAHIonPhase1 = CAHitNtupletCUDAT; +DEFINE_FWK_MODULE(CAHitNtupletCUDAHIonPhase1); diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc index 38ede8eb80d85..3c4125a75b6b1 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc @@ -56,8 +56,7 @@ void CAHitNtupletGeneratorKernelsCPU::buildDoublets(const HitsCon this->device_theCellTracksContainer_); // no need to use the Traits allocations, since we know this is being compiled for the CPU - //this->device_theCells_ = Traits::template make_unique(this->params_.cellCuts_.maxNumberOfDoublets_, stream); - this->device_theCells_ = std::make_unique(this->params_.cellCuts_.maxNumberOfDoublets_); + this->device_theCells_ = std::make_unique(this->params_.caParams_.maxNumberOfDoublets_); if (0 == nhits) return; // protect against empty events @@ -73,7 +72,8 @@ void CAHitNtupletGeneratorKernelsCPU::buildDoublets(const HitsCon hh, this->isOuterHitOfCell_, nActualPairs, - this->params_.cellCuts_); + this->params_.caParams_.maxNumberOfDoublets_, + this->cellCuts_.get()); } template @@ -203,7 +203,7 @@ void CAHitNtupletGeneratorKernelsCPU::classifyTuples(const HitsCo this->device_theCellTracks_.get(), this->isOuterHitOfCell_, nhits, - this->params_.cellCuts_.maxNumberOfDoublets_, + this->params_.caParams_.maxNumberOfDoublets_, this->counters_); } @@ -227,3 +227,4 @@ void CAHitNtupletGeneratorKernelsCPU::classifyTuples(const HitsCo template class CAHitNtupletGeneratorKernelsCPU; template class CAHitNtupletGeneratorKernelsCPU; +template class CAHitNtupletGeneratorKernelsCPU; diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu index 8b3d78f53401a..b13cdcda4d28d 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu @@ -64,7 +64,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(const HitsCon } blockSize = 64; - numberOfBlocks = (3 * this->params_.cellCuts_.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize; + numberOfBlocks = (3 * this->params_.caParams_.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize; kernel_find_ntuplets<<>>(hh, tracks_view, this->device_theCells_.get(), @@ -209,7 +209,7 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(const HitsCon } this->device_theCells_ = - cms::cuda::make_device_unique(this->params_.cellCuts_.maxNumberOfDoublets_, stream); + cms::cuda::make_device_unique(this->params_.caParams_.maxNumberOfDoublets_, stream); #ifdef GPU_DEBUG cudaDeviceSynchronize(); @@ -227,6 +227,7 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(const HitsCon int blocks = (4 * nhits + threadsPerBlock - 1) / threadsPerBlock; dim3 blks(1, blocks, 1); dim3 thrs(stride, threadsPerBlock, 1); + getDoubletsFromHisto<<>>(this->device_theCells_.get(), this->device_nCells_, this->device_theCellNeighbors_.get(), @@ -234,7 +235,8 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(const HitsCon hh, this->isOuterHitOfCell_, nActualPairs, - this->params_.cellCuts_); + this->params_.caParams_.maxNumberOfDoublets_, + this->cellCuts_.get()); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG @@ -329,7 +331,7 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(const HitsCo } if (this->params_.doStats_) { - numberOfBlocks = (std::max(nhits, int(this->params_.cellCuts_.maxNumberOfDoublets_)) + blockSize - 1) / blockSize; + numberOfBlocks = (std::max(nhits, int(this->params_.caParams_.maxNumberOfDoublets_)) + blockSize - 1) / blockSize; kernel_checkOverflows <<>>(tracks_view, this->device_tupleMultiplicity_.get(), @@ -341,7 +343,7 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(const HitsCo this->device_theCellTracks_.get(), this->isOuterHitOfCell_, nhits, - this->params_.cellCuts_.maxNumberOfDoublets_, + this->params_.caParams_.maxNumberOfDoublets_, this->counters_); cudaCheck(cudaGetLastError()); } @@ -387,3 +389,4 @@ void CAHitNtupletGeneratorKernelsGPU::printCounters(Counters cons template class CAHitNtupletGeneratorKernelsGPU; template class CAHitNtupletGeneratorKernelsGPU; +template class CAHitNtupletGeneratorKernelsGPU; diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h index 821826dfef873..55dd7412a46bb 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h @@ -33,6 +33,7 @@ namespace caHitNtupletGenerator { //CAParams struct CACommon { + const uint32_t maxNumberOfDoublets_; const uint32_t minHitsPerNtuplet_; const float ptmin_; const float CAThetaCutBarrel_; @@ -202,6 +203,7 @@ class CAHitNtupletGeneratorKernels { using QualityCuts = pixelTrack::QualityCutsT; using Params = caHitNtupletGenerator::ParamsT; using CAParams = caHitNtupletGenerator::CAParamsT; + using CellCuts = gpuPixelDoublets::CellCutsT; using Counters = caHitNtupletGenerator::Counters; template @@ -226,7 +228,7 @@ class CAHitNtupletGeneratorKernels { using HitContainer = typename TrackSoA::HitContainer; CAHitNtupletGeneratorKernels(Params const& params) - : params_(params), paramsMaxDoubletes3Quarters_(3 * params.cellCuts_.maxNumberOfDoublets_ / 4) {} + : params_(params), paramsMaxDoubletes3Quarters_(3 * params.caParams_.maxNumberOfDoublets_ / 4) {} ~CAHitNtupletGeneratorKernels() = default; @@ -245,6 +247,7 @@ class CAHitNtupletGeneratorKernels { protected: Counters* counters_ = nullptr; + // CellCuts* cellCuts_ = nullptr; // workspace unique_ptr cellStorage_; unique_ptr device_theCellNeighbors_; @@ -261,6 +264,8 @@ class CAHitNtupletGeneratorKernels { unique_ptr device_hitToTupleStorage_; typename HitToTuple::View hitToTupleView_; + unique_ptr cellCuts_; + cms::cuda::AtomicPairCounter* device_hitToTuple_apc_ = nullptr; cms::cuda::AtomicPairCounter* device_hitTuple_apc_ = nullptr; @@ -303,6 +308,8 @@ class CAHitNtupletGeneratorKernelsGPU : public CAHitNtupletGeneratorKernels; using TkSoAView = TrackSoAView; + using Params = caHitNtupletGenerator::ParamsT; + public: void launchKernels(const HitsConstView& hh, TkSoAView& track_view, cudaStream_t cudaStream); void classifyTuples(const HitsConstView& hh, TkSoAView& track_view, cudaStream_t cudaStream); @@ -328,6 +335,8 @@ class CAHitNtupletGeneratorKernelsCPU : public CAHitNtupletGeneratorKernels; using TkSoAView = TrackSoAView; + using Params = caHitNtupletGenerator::ParamsT; + public: void launchKernels(const HitsConstView& hh, TkSoAView& track_view, cudaStream_t cudaStream); void classifyTuples(const HitsConstView& hh, TkSoAView& track_view, cudaStream_t cudaStream); diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc index af085bb12eddd..01cca3187d5f4 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc @@ -2,7 +2,7 @@ #include "CAHitNtupletGeneratorKernels.h" -//#define GPU_DEBUG +// #define GPU_DEBUG template #ifdef __CUDACC__ void CAHitNtupletGeneratorKernelsGPU::allocateOnGPU(int32_t nHits, cudaStream_t stream) { @@ -11,6 +11,9 @@ void CAHitNtupletGeneratorKernelsGPU::allocateOnGPU(int32_t nHits void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU(int32_t nHits, cudaStream_t stream) { using Traits = cms::cudacompat::CPUTraits; #endif + + using CellCuts = gpuPixelDoublets::CellCutsT; + ////////////////////////////////////////////////////////// // ALLOCATIONS FOR THE INTERMEDIATE RESULTS (STAYS ON WORKER) ////////////////////////////////////////////////////////// @@ -38,11 +41,15 @@ void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU(int32_t nHits this->device_hitToTuple_apc_ = (cms::cuda::AtomicPairCounter*)this->device_storage_.get() + 1; this->device_nCells_ = (uint32_t*)(this->device_storage_.get() + 2); + this->cellCuts_ = Traits::template make_unique(stream); // FIXME: consider collapsing these 3 in one adhoc kernel if constexpr (std::is_same::value) { cudaCheck(cudaMemsetAsync(this->device_nCells_, 0, sizeof(uint32_t), stream)); + cudaCheck(cudaMemcpyAsync( + this->cellCuts_.get(), &(this->params_.cellCuts_), sizeof(CellCuts), cudaMemcpyDefault, stream)); } else { *(this->device_nCells_) = 0; + *(this->cellCuts_.get()) = this->params_.cellCuts_; } cms::cuda::launchZero(this->device_tupleMultiplicity_.get(), stream); cms::cuda::launchZero(this->hitToTupleView_, stream); // we may wish to keep it in the edm @@ -54,6 +61,8 @@ void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU(int32_t nHits template class CAHitNtupletGeneratorKernelsGPU; template class CAHitNtupletGeneratorKernelsGPU; +template class CAHitNtupletGeneratorKernelsGPU; template class CAHitNtupletGeneratorKernelsCPU; template class CAHitNtupletGeneratorKernelsCPU; +template class CAHitNtupletGeneratorKernelsCPU; diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h index 85386305eca6a..540c0b92f9015 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsImpl.h @@ -108,7 +108,7 @@ namespace caHitNtupletGeneratorKernels { for (int idx = first, nt = tracks_view.hitIndices().nOnes(); idx < nt; idx += gridDim.x * blockDim.x) { if (tracks_view.hitIndices().size(idx) > TrackerTraits::maxHitsOnTrack) // current real limit printf("ERROR %d, %d\n", idx, tracks_view.hitIndices().size(idx)); - assert(ftracks_view.hitIndices().size(idx) <= TrackerTraits::maxHitsOnTrack); + assert(tracks_view.hitIndices().size(idx) <= TrackerTraits::maxHitsOnTrack); for (auto ih = tracks_view.hitIndices().begin(idx); ih != tracks_view.hitIndices().end(idx); ++ih) assert(int(*ih) < nHits); } @@ -130,17 +130,7 @@ namespace caHitNtupletGeneratorKernels { cellNeighbors->size(), cellTracks->size(), hitToTuple->size()); -// printf("cellTracksSizes;"); -// for (int i = 0; i < cellTracks->size(); i++) { -// printf("%d;",cellTracks[i].size()); -// } -// -// printf("\n"); -// printf("cellNeighborsSizes;"); -// for (int i = 0; i < cellNeighbors->size(); i++) { -// printf("%d;",cellNeighbors[i].size()); -// } -// printf("\n"); + #endif } diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc index 6765703f35a73..5e305bebd0e95 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc @@ -61,13 +61,16 @@ namespace { template struct topologyCuts> { static constexpr CAParamsT makeCACuts(edm::ParameterSet const& cfg) { - return CAParamsT{{cfg.getParameter("minHitsPerNtuplet"), - (float)cfg.getParameter("ptmin"), - (float)cfg.getParameter("CAThetaCutBarrel"), - (float)cfg.getParameter("CAThetaCutForward"), - (float)cfg.getParameter("hardCurvCut"), - (float)cfg.getParameter("dcaCutInnerTriplet"), - (float)cfg.getParameter("dcaCutOuterTriplet")}}; + return CAParamsT{{ + cfg.getParameter("maxNumberOfDoublets"), + cfg.getParameter("minHitsPerNtuplet"), + (float)cfg.getParameter("ptmin"), + (float)cfg.getParameter("CAThetaCutBarrel"), + (float)cfg.getParameter("CAThetaCutForward"), + (float)cfg.getParameter("hardCurvCut"), + (float)cfg.getParameter("dcaCutInnerTriplet"), + (float)cfg.getParameter("dcaCutOuterTriplet"), + }}; }; static constexpr pixelTrack::QualityCutsT makeQualityCuts(edm::ParameterSet const& pset) { @@ -95,7 +98,8 @@ namespace { template struct topologyCuts> { static constexpr CAParamsT makeCACuts(edm::ParameterSet const& cfg) { - return CAParamsT{{cfg.getParameter("minHitsPerNtuplet"), + return CAParamsT{{cfg.getParameter("maxNumberOfDoublets"), + cfg.getParameter("minHitsPerNtuplet"), (float)cfg.getParameter("ptmin"), (float)cfg.getParameter("CAThetaCutBarrel"), (float)cfg.getParameter("CAThetaCutForward"), @@ -118,13 +122,15 @@ namespace { //Cell Cuts, as they are the cuts have the same logic for Phase2 and Phase1 //keeping them separate would allow further differentiation in the future //moving them to topologyCuts and using the same syntax - template - CellCutsT makeCellCuts(edm::ParameterSet const& cfg) { - return CellCutsT{cfg.getParameter("maxNumberOfDoublets"), - cfg.getParameter("doClusterCut"), + template + CellCutsT makeCellCuts(edm::ParameterSet const& cfg) { + return CellCutsT{cfg.getParameter("doClusterCut"), cfg.getParameter("doZ0Cut"), cfg.getParameter("doPtCut"), - cfg.getParameter("idealConditions")}; + cfg.getParameter("idealConditions"), + (float)cfg.getParameter("z0Cut"), + (float)cfg.getParameter("ptCut"), + cfg.getParameter>("phiCuts")}; } } // namespace @@ -174,6 +180,8 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::Pa desc.add("idealConditions", true); desc.add("includeJumpingForwardDoublets", false); + desc.add("z0Cut", 12.0f); + desc.add("ptCut", 0.5f); edm::ParameterSetDescription trackQualityCuts; trackQualityCuts.add("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut"); @@ -188,6 +196,44 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::Pa trackQualityCuts.add("quadrupletMinPt", 0.3)->setComment("Min pT for quadruplets, in GeV"); trackQualityCuts.add("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm"); trackQualityCuts.add("quadrupletMaxZip", 12.)->setComment("Max |Zip| for quadruplets, in cm"); + + desc.add>("phiCuts", + std::vector(phase1PixelTopology::phicuts, std::end(phase1PixelTopology::phicuts))) + ->setComment("Cuts in phi for cells"); + + desc.add("trackQualityCuts", trackQualityCuts) + ->setComment( + "Quality cuts based on the results of the track fit:\n - apply a pT-dependent chi2 cut;\n - apply \"region " + "cuts\" based on the fit results (pT, Tip, Zip)."); +} + +template <> +void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::ParameterSetDescription& desc) { + fillDescriptionsCommon(desc); + + desc.add("idealConditions", false); + desc.add("includeJumpingForwardDoublets", false); + desc.add("z0Cut", 10.0f); + desc.add("ptCut", 0.0f); + + edm::ParameterSetDescription trackQualityCuts; + trackQualityCuts.add("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut"); + trackQualityCuts.add>("chi2Coeff", {0.9, 1.8})->setComment("chi2 at 1GeV and at ptMax above"); + trackQualityCuts.add("chi2Scale", 8.) + ->setComment( + "Factor to multiply the pT-dependent chi2 cut (currently: 8 for the broken line fit, ?? for the Riemann " + "fit)"); + trackQualityCuts.add("tripletMinPt", 0.0)->setComment("Min pT for triplets, in GeV"); + trackQualityCuts.add("tripletMaxTip", 0.1)->setComment("Max |Tip| for triplets, in cm"); + trackQualityCuts.add("tripletMaxZip", 6.)->setComment("Max |Zip| for triplets, in cm"); + trackQualityCuts.add("quadrupletMinPt", 0.0)->setComment("Min pT for quadruplets, in GeV"); + trackQualityCuts.add("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm"); + trackQualityCuts.add("quadrupletMaxZip", 6.)->setComment("Max |Zip| for quadruplets, in cm"); + + desc.add>("phiCuts", + std::vector(phase1PixelTopology::phicuts, std::end(phase1PixelTopology::phicuts))) + ->setComment("Cuts in phi for cells"); + desc.add("trackQualityCuts", trackQualityCuts) ->setComment( "Quality cuts based on the results of the track fit:\n - apply a pT-dependent chi2 cut;\n - apply \"region " @@ -201,12 +247,19 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::Pa desc.add("idealConditions", false); desc.add("includeFarForwards", true); desc.add("includeJumpingForwardDoublets", true); + desc.add("z0Cut", 7.5f); + desc.add("ptCut", 0.85f); edm::ParameterSetDescription trackQualityCuts; trackQualityCuts.add("maxChi2", 5.)->setComment("Max normalized chi2"); trackQualityCuts.add("minPt", 0.5)->setComment("Min pT in GeV"); trackQualityCuts.add("maxTip", 0.3)->setComment("Max |Tip| in cm"); trackQualityCuts.add("maxZip", 12.)->setComment("Max |Zip|, in cm"); + + desc.add>("phiCuts", + std::vector(phase2PixelTopology::phicuts, std::end(phase2PixelTopology::phicuts))) + ->setComment("Cuts in phi for cells"); + desc.add("trackQualityCuts", trackQualityCuts) ->setComment( "Quality cuts based on the results of the track fit:\n - apply cuts based on the fit results (pT, Tip, " @@ -360,3 +413,4 @@ TrackSoAHeterogeneousHost CAHitNtupletGeneratorOnGPU; template class CAHitNtupletGeneratorOnGPU; +template class CAHitNtupletGeneratorOnGPU; diff --git a/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc b/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc index befd30ffab7b2..c36ed924911f0 100644 --- a/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc +++ b/RecoTracker/PixelSeeding/plugins/HelixFitOnGPU.cc @@ -19,3 +19,4 @@ void HelixFitOnGPU::deallocateOnGPU() {} template class HelixFitOnGPU; template class HelixFitOnGPU; +template class HelixFitOnGPU; diff --git a/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cc b/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cc index 2678f60f75b3f..4f9037da0ceb7 100644 --- a/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cc +++ b/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cc @@ -117,3 +117,4 @@ void HelixFitOnGPU::launchRiemannKernelsOnCPU(const TrackingRecHi template class HelixFitOnGPU; template class HelixFitOnGPU; +template class HelixFitOnGPU; diff --git a/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cu b/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cu index 99c55992bbf71..73b33cc336fce 100644 --- a/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cu +++ b/RecoTracker/PixelSeeding/plugins/RiemannFitOnGPU.cu @@ -133,3 +133,4 @@ void HelixFitOnGPU::launchRiemannKernels(const TrackingRecHitSoAC template class HelixFitOnGPU; template class HelixFitOnGPU; +template class HelixFitOnGPU; diff --git a/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h b/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h index 13e32cce0bfa5..7e97bc29b04d1 100644 --- a/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h +++ b/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h @@ -46,10 +46,11 @@ namespace gpuPixelDoublets { HitsConstView hh, OuterHitOfCell isOuterHitOfCell, int nActualPairs, - CellCutsT cuts) { + const int maxNumOfDoublets, + CellCutsT* const cuts) { doubletsFromHisto( - nActualPairs, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, cuts); + nActualPairs, maxNumOfDoublets, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, cuts); } } // namespace gpuPixelDoublets diff --git a/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h b/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h index 3d9a054f68eda..c6f6dcfa0f822 100644 --- a/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h +++ b/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h @@ -16,8 +16,8 @@ #include "CAStructures.h" #include "GPUCACell.h" -//#define GPU_DEBUG -//#define NTUPLE_DEBUG +// #define GPU_DEBUG +// #define NTUPLE_DEBUG namespace gpuPixelDoublets { @@ -39,11 +39,33 @@ namespace gpuPixelDoublets { using H = HitsConstView; using T = TrackerTraits; - const uint32_t maxNumberOfDoublets_; - const bool doClusterCut_; - const bool doZ0Cut_; - const bool doPtCut_; - const bool idealConditions_; //this is actually not used by phase2 + CellCutsT() = default; + CellCutsT(const bool doClusterCut, + const bool doZ0Cut, + const bool doPtCut, + const bool idealConditions, + const float z0Cut, + const float ptCut, + const std::vector& phiCutsV) + : doClusterCut_(doClusterCut), + doZ0Cut_(doZ0Cut), + doPtCut_(doPtCut), + idealConditions_(idealConditions), + z0Cut_(z0Cut), + ptCut_(ptCut) { + assert(phiCutsV.size() == TrackerTraits::nPairs); + std::copy(phiCutsV.begin(), phiCutsV.end(), &phiCuts[0]); + } + + bool doClusterCut_; + bool doZ0Cut_; + bool doPtCut_; + bool idealConditions_; //this is actually not used by phase2 + + float z0Cut_; + float ptCut_; + + int phiCuts[T::nPairs]; __device__ __forceinline__ bool zSizeCut(H hh, int i, int o) const { const uint32_t mi = hh[i].detectorIndex(); @@ -97,20 +119,26 @@ namespace gpuPixelDoublets { template __device__ __forceinline__ void doubletsFromHisto(uint32_t nPairs, + uint32_t maxNumOfDoublets, GPUCACellT* cells, uint32_t* nCells, CellNeighborsVector* cellNeighbors, CellTracksVector* cellTracks, HitsConstView hh, OuterHitOfCell isOuterHitOfCell, - CellCutsT const& cuts) { + CellCutsT const* cuts) { // ysize cuts (z in the barrel) times 8 // these are used if doClusterCut is true - const bool doClusterCut = cuts.doClusterCut_; - const bool doZ0Cut = cuts.doZ0Cut_; - const bool doPtCut = cuts.doPtCut_; - const uint32_t maxNumOfDoublets = cuts.maxNumberOfDoublets_; + const bool doClusterCut = cuts->doClusterCut_; + const bool doZ0Cut = cuts->doZ0Cut_; + const bool doPtCut = cuts->doPtCut_; + + const float z0cut = cuts->z0Cut_; // cm + const float hardPtCut = cuts->ptCut_; // GeV + // cm (1 GeV track has 1 GeV/c / (e * 3.8T) ~ 87 cm radius in a 3.8T field) + const float minRadius = hardPtCut * 87.78f; + const float minRadius2T4 = 4.f * minRadius * minRadius; using PhiBinner = typename TrackingRecHitSoA::PhiBinner; @@ -178,18 +206,12 @@ namespace gpuPixelDoublets { if (mez < TrackerTraits::minz[pairLayerId] || mez > TrackerTraits::maxz[pairLayerId]) continue; - if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts.clusterCut(hh, i)) + if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts->clusterCut(hh, i)) continue; auto mep = hh[i].iphi(); auto mer = hh[i].rGlobal(); - // all cuts: true if fails - constexpr float z0cut = TrackerTraits::z0Cut; // cm - constexpr float hardPtCut = TrackerTraits::doubletHardPt; // GeV - // cm (1 GeV track has 1 GeV/c / (e * 3.8T) ~ 87 cm radius in a 3.8T field) - constexpr float minRadius = hardPtCut * 87.78f; - constexpr float minRadius2T4 = 4.f * minRadius * minRadius; auto ptcut = [&](int j, int16_t idphi) { auto r2t4 = minRadius2T4; auto ri = mer; @@ -204,7 +226,7 @@ namespace gpuPixelDoublets { return dr > TrackerTraits::maxr[pairLayerId] || dr < 0 || std::abs((mez * ro - mer * zo)) > z0cut * dr; }; - auto iphicut = TrackerTraits::phicuts[pairLayerId]; + auto iphicut = cuts->phiCuts[pairLayerId]; auto kl = PhiBinner::bin(int16_t(mep - iphicut)); auto kh = PhiBinner::bin(int16_t(mep + iphicut)); @@ -242,7 +264,7 @@ namespace gpuPixelDoublets { if (idphi > iphicut) continue; - if (doClusterCut && cuts.zSizeCut(hh, i, oi)) + if (doClusterCut && cuts->zSizeCut(hh, i, oi)) continue; if (doPtCut && ptcut(oi, idphi)) continue; diff --git a/RecoTracker/PixelSeeding/test/trip_cfg.py b/RecoTracker/PixelSeeding/test/trip_cfg.py index a0603b9736c06..46d2fa36b8ca8 100644 --- a/RecoTracker/PixelSeeding/test/trip_cfg.py +++ b/RecoTracker/PixelSeeding/test/trip_cfg.py @@ -44,8 +44,8 @@ # cout = cms.untracked.PSet( threshold = cms.untracked.string('INFO')) #) -from RecoPixelVertexing.PixelTriplets.PixelTripletHLTGenerator_cfi import * -from RecoPixelVertexing.PixelTriplets.PixelTripletLargeTipGenerator_cfi import * +from RecoTracker.PixelSeeding.PixelTripletHLTGenerator_cfi import * +from RecoTracker.PixelSeeding.PixelTripletLargeTipGenerator_cfi import * from RecoTracker.TkTrackingRegions.GlobalTrackingRegion_cfi import * process.triplets = cms.EDAnalyzer("HitTripletProducer", diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc index 7245d7f9eb7e5..6bff9a7c42292 100644 --- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc +++ b/RecoTracker/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc @@ -99,3 +99,6 @@ DEFINE_FWK_MODULE(PixelTrackDumpCUDAPhase1); using PixelTrackDumpCUDAPhase2 = PixelTrackDumpCUDAT; DEFINE_FWK_MODULE(PixelTrackDumpCUDAPhase2); + +using PixelTrackDumpCUDAHIonPhase1 = PixelTrackDumpCUDAT; +DEFINE_FWK_MODULE(PixelTrackDumpCUDAHIonPhase1); diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index 465f210b0cb0e..fe65af40b4dd5 100644 --- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc +++ b/RecoTracker/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc @@ -260,3 +260,6 @@ DEFINE_FWK_MODULE(PixelTrackProducerFromSoAPhase1); using PixelTrackProducerFromSoAPhase2 = PixelTrackProducerFromSoAT; DEFINE_FWK_MODULE(PixelTrackProducerFromSoAPhase2); + +using PixelTrackProducerFromSoAHIonPhase1 = PixelTrackProducerFromSoAT; +DEFINE_FWK_MODULE(PixelTrackProducerFromSoAHIonPhase1); diff --git a/RecoTracker/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc b/RecoTracker/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc index 8f8ca96833cbb..fc2c76ff00155 100644 --- a/RecoTracker/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc +++ b/RecoTracker/PixelTrackFitting/plugins/PixelTrackSoAFromCUDA.cc @@ -108,3 +108,6 @@ DEFINE_FWK_MODULE(PixelTrackSoAFromCUDAPhase1); using PixelTrackSoAFromCUDAPhase2 = PixelTrackSoAFromCUDAT; DEFINE_FWK_MODULE(PixelTrackSoAFromCUDAPhase2); + +using PixelTrackSoAFromCUDAHIonPhase1 = PixelTrackSoAFromCUDAT; +DEFINE_FWK_MODULE(PixelTrackSoAFromCUDAHIonPhase1); diff --git a/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py b/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py index 05340f69ea2fe..35051e7b50cf6 100644 --- a/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py +++ b/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py @@ -93,8 +93,12 @@ from RecoTracker.PixelSeeding.caHitNtupletCUDAPhase1_cfi import caHitNtupletCUDAPhase1 as _pixelTracksCUDA from RecoTracker.PixelSeeding.caHitNtupletCUDAPhase2_cfi import caHitNtupletCUDAPhase2 as _pixelTracksCUDAPhase2 +from RecoTracker.PixelSeeding.caHitNtupletCUDAHIonPhase1_cfi import caHitNtupletCUDAHIonPhase1 as _pixelTracksCUDAHIonPhase1 +# Phase 2 modifier from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker +# HIon modifiers +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA # SwitchProducer providing the pixel tracks in SoA format on the CPU pixelTracksSoA = SwitchProducerCUDA( @@ -114,6 +118,7 @@ # convert the pixel tracks from SoA to legacy format from RecoTracker.PixelTrackFitting.pixelTrackProducerFromSoAPhase1_cfi import pixelTrackProducerFromSoAPhase1 as _pixelTrackProducerFromSoA from RecoTracker.PixelTrackFitting.pixelTrackProducerFromSoAPhase2_cfi import pixelTrackProducerFromSoAPhase2 as _pixelTrackProducerFromSoAPhase2 +from RecoTracker.PixelTrackFitting.pixelTrackProducerFromSoAHIonPhase1_cfi import pixelTrackProducerFromSoAHIonPhase1 as _pixelTrackProducerFromSoAHIonPhase1 pixelNtupletFit.toReplaceWith(pixelTracks, _pixelTrackProducerFromSoA.clone( pixelRecHitLegacySrc = "siPixelRecHitsPreSplitting", @@ -123,6 +128,10 @@ pixelRecHitLegacySrc = "siPixelRecHitsPreSplitting", )) +(pixelNtupletFit & pp_on_AA).toReplaceWith(pixelTracks, _pixelTrackProducerFromSoAHIonPhase1.clone( + pixelRecHitLegacySrc = "siPixelRecHitsPreSplitting", +)) + pixelNtupletFit.toReplaceWith(pixelTracksTask, cms.Task( # build the pixel ntuplets and the pixel tracks in SoA format on the GPU pixelTracksSoA, @@ -148,6 +157,7 @@ # SwitchProducer providing the pixel tracks in SoA format on the CPU from RecoTracker.PixelTrackFitting.pixelTrackSoAFromCUDAPhase1_cfi import pixelTrackSoAFromCUDAPhase1 as _pixelTracksSoA from RecoTracker.PixelTrackFitting.pixelTrackSoAFromCUDAPhase2_cfi import pixelTrackSoAFromCUDAPhase2 as _pixelTracksSoAPhase2 +from RecoTracker.PixelTrackFitting.pixelTrackSoAFromCUDAHIonPhase1_cfi import pixelTrackSoAFromCUDAHIonPhase1 as _pixelTracksSoAHIonPhase1 gpu.toModify(pixelTracksSoA, # transfer the pixel tracks in SoA format to the host @@ -157,16 +167,29 @@ (gpu & phase2_tracker).toModify(pixelTracksSoA,cuda = _pixelTracksSoAPhase2.clone( )) +(gpu & pp_on_AA).toModify(pixelTracksSoA,cuda = _pixelTracksSoAHIonPhase1.clone( +)) + phase2_tracker.toModify(pixelTracksSoA,cpu = _pixelTracksCUDAPhase2.clone( pixelRecHitSrc = "siPixelRecHitsPreSplittingSoA", onGPU = False )) +pp_on_AA.toModify(pixelTracksSoA,cpu = _pixelTracksCUDAHIonPhase1.clone( + pixelRecHitSrc = "siPixelRecHitsPreSplittingSoA", + onGPU = False +)) + phase2_tracker.toReplaceWith(pixelTracksCUDA,_pixelTracksCUDAPhase2.clone( pixelRecHitSrc = "siPixelRecHitsPreSplittingCUDA", onGPU = True, )) +pp_on_AA.toReplaceWith(pixelTracksCUDA,_pixelTracksCUDAHIonPhase1.clone( + pixelRecHitSrc = "siPixelRecHitsPreSplittingCUDA", + onGPU = True, +)) + (pixelNtupletFit & gpu).toReplaceWith(pixelTracksTask, cms.Task( # build the pixel ntuplets and pixel tracks in SoA format on the GPU pixelTracksCUDA, diff --git a/RecoTracker/PixelVertexFinding/plugins/PixelVertexProducerCUDA.cc b/RecoTracker/PixelVertexFinding/plugins/PixelVertexProducerCUDA.cc index 0c4519a3d3724..a1f4101252319 100644 --- a/RecoTracker/PixelVertexFinding/plugins/PixelVertexProducerCUDA.cc +++ b/RecoTracker/PixelVertexFinding/plugins/PixelVertexProducerCUDA.cc @@ -65,6 +65,7 @@ PixelVertexProducerCUDAT::PixelVertexProducerCUDAT(const edm::Par conf.getParameter("useDensity"), conf.getParameter("useDBSCAN"), conf.getParameter("useIterative"), + conf.getParameter("doSplitting"), conf.getParameter("minT"), conf.getParameter("eps"), conf.getParameter("errmax"), @@ -92,6 +93,7 @@ void PixelVertexProducerCUDAT::fillDescriptions(edm::Configuratio desc.add("useDensity", true); desc.add("useDBSCAN", false); desc.add("useIterative", false); + desc.add("doSplitting", true); desc.add("minT", 2); // min number of neighbours to be "core" desc.add("eps", 0.07); // max absolute distance to cluster @@ -159,3 +161,6 @@ DEFINE_FWK_MODULE(PixelVertexProducerCUDAPhase1); using PixelVertexProducerCUDAPhase2 = PixelVertexProducerCUDAT; DEFINE_FWK_MODULE(PixelVertexProducerCUDAPhase2); + +using PixelVertexProducerCUDAHIonPhase1 = PixelVertexProducerCUDAT; +DEFINE_FWK_MODULE(PixelVertexProducerCUDAHIonPhase1); diff --git a/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.cc b/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.cc index 950a31f8ac48a..73fa1408a0aab 100644 --- a/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.cc +++ b/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.cc @@ -169,11 +169,13 @@ namespace gpuVertexFinder { cudaCheck(cudaGetLastError()); fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFirstFit); cudaCheck(cudaGetLastError()); - // one block per vertex... - splitVerticesKernel<<>>(soa, ws_d.view(), maxChi2ForSplit); - cudaCheck(cudaGetLastError()); - fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFinalFit); - cudaCheck(cudaGetLastError()); + if (doSplitting_) { + // one block per vertex... + splitVerticesKernel<<>>(soa, ws_d.view(), maxChi2ForSplit); + cudaCheck(cudaGetLastError()); + fitVerticesKernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view(), maxChi2ForFinalFit); + cudaCheck(cudaGetLastError()); + } sortByPt2Kernel<<<1, maxThreadsForPrint, 0, stream>>>(soa, ws_d.view()); } cudaCheck(cudaGetLastError()); @@ -190,8 +192,10 @@ namespace gpuVertexFinder { #endif // PIXVERTEX_DEBUG_PRODUCE fitVertices(soa, ws_d.view(), maxChi2ForFirstFit); // one block per vertex! - splitVertices(soa, ws_d.view(), maxChi2ForSplit); - fitVertices(soa, ws_d.view(), maxChi2ForFinalFit); + if (doSplitting_) { + splitVertices(soa, ws_d.view(), maxChi2ForSplit); + fitVertices(soa, ws_d.view(), maxChi2ForFinalFit); + } sortByPt2(soa, ws_d.view()); #endif @@ -200,4 +204,5 @@ namespace gpuVertexFinder { template class Producer; template class Producer; + template class Producer; } // namespace gpuVertexFinder diff --git a/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.h b/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.h index 3839424b56530..59ec304b67521 100644 --- a/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.h +++ b/RecoTracker/PixelVertexFinding/plugins/gpuVertexFinder.h @@ -31,6 +31,7 @@ namespace gpuVertexFinder { bool useDensity, bool useDBSCAN, bool useIterative, + bool doSplitting, int iminT, // min number of neighbours to be "core" float ieps, // max absolute distance to cluster float ierrmax, // max error to be "seed" @@ -40,6 +41,7 @@ namespace gpuVertexFinder { useDensity_(useDensity), useDBSCAN_(useDBSCAN), useIterative_(useIterative), + doSplitting_(doSplitting), minT(iminT), eps(ieps), errmax(ierrmax), @@ -55,6 +57,7 @@ namespace gpuVertexFinder { const bool useDensity_; const bool useDBSCAN_; const bool useIterative_; + const bool doSplitting_; int minT; // min number of neighbours to be "core" float eps; // max absolute distance to cluster diff --git a/Validation/RecoTrack/python/PostProcessorTracker_cfi.py b/Validation/RecoTrack/python/PostProcessorTracker_cfi.py index 54fa0364fe239..ac9686b547701 100644 --- a/Validation/RecoTrack/python/PostProcessorTracker_cfi.py +++ b/Validation/RecoTrack/python/PostProcessorTracker_cfi.py @@ -230,6 +230,7 @@ def _addNoFlow(module): noFlowDists = cms.untracked.vstring(), outputFileName = cms.untracked.string("") ) + _addNoFlow(postProcessorTrack) postProcessorTrack2D = DQMEDHarvester("DQMGenericClient", @@ -246,6 +247,7 @@ def _addNoFlow(module): noFlowDists = cms.untracked.vstring(), outputFileName = cms.untracked.string("") ) + _addNoFlow(postProcessorTrack2D) # nrec/nsim makes sense only for @@ -339,7 +341,6 @@ def _addNoFlow(module): phase2_tracker.toReplaceWith(postProcessorTrack,postProcessorTrackPhase2) phase2_tracker.toReplaceWith(postProcessorTrackSummary,postProcessorTrackSummaryPhase2) - from Configuration.ProcessModifiers.displacedTrackValidation_cff import displacedTrackValidation postProcessorTrackDisplaced = postProcessorTrack.clone() postProcessorTrackDisplaced.subDirs.extend(["Tracking/TrackDisplaced/*"]) @@ -348,11 +349,20 @@ def _addNoFlow(module): displacedTrackValidation.toReplaceWith(postProcessorTrack,postProcessorTrackDisplaced) displacedTrackValidation.toReplaceWith(postProcessorTrackSummary,postProcessorTrackSummaryDisplaced) +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA + +_defaultSubdirsHIon = _defaultSubdirs + ["Tracking/HIPixelTrack/*"] +_defaultSubdirsSummaryHIon = [e.replace("/*","") for e in _defaultSubdirsHIon] + +pp_on_AA.toModify(postProcessorTrack,subDirs = _defaultSubdirsHIon) +pp_on_AA.toModify(postProcessorTrack2D,subDirs = _defaultSubdirsHIon) +pp_on_AA.toModify(postProcessorTrackSummary,subDirs = _defaultSubdirsHIon) + postProcessorTrackTrackingOnly = postProcessorTrack.clone() -postProcessorTrackTrackingOnly.subDirs.extend(["Tracking/TrackBHadron/*", "Tracking/TrackSeeding/*", "Tracking/PixelTrack/*", "Tracking/PixelTrackFromPV/*", "Tracking/PixelTrackFromPVAllTP/*", "Tracking/PixelTrackBHadron/*"]) +postProcessorTrackTrackingOnly.subDirs.extend(["Tracking/TrackBHadron/*", "Tracking/TrackSeeding/*", "Tracking/HIPixelTrack/*", "Tracking/PixelTrack/*", "Tracking/PixelTrackFromPV/*", "Tracking/PixelTrackFromPVAllTP/*", "Tracking/PixelTrackBHadron/*"]) postProcessorTrackSummaryTrackingOnly = postProcessorTrackSummary.clone() -postProcessorTrackSummaryTrackingOnly.subDirs.extend(["Tracking/TrackBHadron", "Tracking/TrackSeeding", "Tracking/PixelTrack", "Tracking/PixelTrackFromPV", "Tracking/PixelTrackFromPVAllTP", "Tracking/PixelTrackBHadron"]) +postProcessorTrackSummaryTrackingOnly.subDirs.extend(["Tracking/TrackBHadron", "Tracking/TrackSeeding", "Tracking/HIPixelTrack", "Tracking/PixelTrack", "Tracking/PixelTrackFromPV", "Tracking/PixelTrackFromPVAllTP", "Tracking/PixelTrackBHadron"]) postProcessorTrackSequenceTrackingOnly = cms.Sequence( postProcessorTrackTrackingOnly+ diff --git a/Validation/RecoTrack/python/TrackValidation_cff.py b/Validation/RecoTrack/python/TrackValidation_cff.py index 4976cd18e321d..02738c0c40a35 100644 --- a/Validation/RecoTrack/python/TrackValidation_cff.py +++ b/Validation/RecoTrack/python/TrackValidation_cff.py @@ -711,6 +711,29 @@ def _uniqueFirstLayers(layerList): VertexAssociatorByPositionAndTracks, trackingParticleNumberOfLayersProducer ) + +# HIon modifiers +from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA + +trackingParticleHIPixelTrackAssociation = trackingParticleRecoTrackAsssociation.clone( + label_tr = "hiConformalPixelTracks", + associator = "quickTrackAssociatorByHits", +) + +from Configuration.ProcessModifiers.pixelNtupletFit_cff import pixelNtupletFit + +pixelNtupletFit.toModify(trackingParticleHIPixelTrackAssociation, + associator = "quickTrackAssociatorByHitsPreSplitting") + +HIPixelVertexAssociatorByPositionAndTracks = VertexAssociatorByPositionAndTracks.clone( + trackAssociation = "trackingParticleHIPixelTrackAssociation" +) + +pp_on_AA.toReplaceWith(tracksValidationTruth, cms.Task( + tracksValidationTruth.copy(), + trackingParticleHIPixelTrackAssociation, + HIPixelVertexAssociatorByPositionAndTracks +)) fastSim.toModify(tracksValidationTruth, lambda x: x.remove(tpClusterProducer)) tracksPreValidation = cms.Task( @@ -745,6 +768,29 @@ def _uniqueFirstLayers(layerList): tracksPreValidation ) +trackValidatorHILowPtConformalValidator = trackValidator.clone( + dirName = "Tracking/HIPixelTrack/", + label = [ + "hiConformalPixelTracks", + ], + doResolutionPlotsForLabels = ["hiConformalPixelTracks"], + trackCollectionForDrCalculation = "hiConformalPixelTracks", + associators = ["trackingParticleHIPixelTrackAssociation"], + vertexAssociator = "HIPixelVertexAssociatorByPositionAndTracks", + dodEdxPlots = False, + cores = cms.InputTag(""), +) + +tracksValidationHIonTask = cms.Task(trackValidatorHILowPtConformalValidator) + +tracksValidationHIon = cms.Sequence( + tracksValidation.copy(), + tracksValidationHIonTask +) + +pp_on_AA.toReplaceWith(tracksValidation,tracksValidationHIon) + + from Configuration.ProcessModifiers.seedingDeepCore_cff import seedingDeepCore seedingDeepCore.toReplaceWith(tracksValidation, cms.Sequence(tracksValidation.copy()+trackValidatorJetCore)) @@ -977,6 +1023,15 @@ def _uniqueFirstLayers(layerList): tracksValidationSeedSelectorsTrackingOnly ) + +tracksValidationHIonTrackingOnly = cms.Sequence( + tracksValidation.copy(), + tracksValidationHIonTask +) + +pp_on_AA.toReplaceWith(tracksValidationTrackingOnly,tracksValidationHIonTrackingOnly) + + #################################################################################################### ### Pixel tracking only mode (placeholder for now) trackingParticlePixelTrackAsssociation = trackingParticleRecoTrackAsssociation.clone( @@ -1035,8 +1090,9 @@ def _uniqueFirstLayers(layerList): label_vertex = "pixelVertices", vertexAssociator = "PixelVertexAssociatorByPositionAndTracks", dodEdxPlots = False, - cores = cms.InputTag(""), + cores = cms.InputTag("") ) + trackValidatorFromPVPixelTrackingOnly = trackValidatorPixelTrackingOnly.clone( dirName = "Tracking/PixelTrackFromPV/", label = [ diff --git a/Validation/RecoTrack/python/plotting/trackingPlots.py b/Validation/RecoTrack/python/plotting/trackingPlots.py index b5b34804d1ec5..b88dac7bcbf81 100644 --- a/Validation/RecoTrack/python/plotting/trackingPlots.py +++ b/Validation/RecoTrack/python/plotting/trackingPlots.py @@ -1378,7 +1378,7 @@ def _appendPixelTrackingPlots(lastDirName, name): _appendPixelTrackingPlots("PixelTrackFromPV", "pixelFromPV") _appendPixelTrackingPlots("PixelTrackFromPVAllTP", "pixelFromPVAllTP") _appendPixelTrackingPlots("PixelTrackBHadron", "pixelbhadron") - +_appendPixelTrackingPlots("HIPixelTrack", "hiPixel") # MiniAOD plotter.append("packedCandidate", _trackingFolders("PackedCandidate"), From 0c04805cf8796429f3a0fba7b134d5d13f5727fb Mon Sep 17 00:00:00 2001 From: AdrianoDee Date: Mon, 10 Jul 2023 16:44:23 +0200 Subject: [PATCH 2/2] Addressing comments --- .../interface/gpuClusteringConstants.h | 7 +-- .../interface/TrackingRecHitSoADevice.h | 1 - .../interface/TrackingRecHitSoAHost.h | 1 - .../PyReleaseValidation/python/relval_gpu.py | 3 +- .../python/upgradeWorkflowComponents.py | 37 ++-------------- .../plugins/SiPixelCompareRecHitsSoA.cc | 1 - .../plugins/SiPixelCompareTrackSoA.cc | 1 - .../plugins/SiPixelMonitorRecHitsSoA.cc | 1 - .../SiPixelHeterogenousDQM_FirstStep_cff.py | 2 +- .../interface/SimplePixelTopology.h | 5 +++ .../plugins/SiPixelDigisClustersFromSoA.cc | 4 +- .../plugins/SiPixelPhase2DigiToClusterCUDA.cc | 8 ++-- .../plugins/SiPixelRawToClusterCUDA.cc | 13 +++--- .../plugins/gpuClusterChargeCut.h | 17 +++---- .../python/siPixelClustersPreSplitting_cff.py | 10 ++--- .../SiPixelClusterizer/test/gpuClustering_t.h | 4 +- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 2 +- .../python/SiPixelRecHits_cfi.py | 11 +++-- .../python/RecoPixelVertexing_cff.py | 4 +- .../plugins/CAHitNtupletGeneratorKernels.cc | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cu | 2 +- .../plugins/CAHitNtupletGeneratorKernels.h | 4 +- .../CAHitNtupletGeneratorKernelsAlloc.cc | 6 +-- .../plugins/CAHitNtupletGeneratorOnGPU.cc | 44 +++++++++---------- .../PixelSeeding/plugins/gpuPixelDoublets.h | 2 +- .../plugins/gpuPixelDoubletsAlgos.h | 19 ++++---- .../python/PixelTracks_cff.py | 4 +- .../python/PostProcessorTracker_cfi.py | 20 ++++++--- 28 files changed, 106 insertions(+), 129 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index cb2b7ace8bd17..923ebaaa5446c 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -13,12 +13,13 @@ namespace gpuClustering { // tested on MC events with 55-75 pileup events constexpr uint32_t maxHitsInIter() { return 160; } //TODO better tuning for PU 140-200 #endif - constexpr uint32_t maxHitsInModule() { return 2048; } - constexpr uint32_t maxNumDigis = 3 * 256 * 1024; // @PU=200 µ=530 sigma=50k this is >4sigma away + constexpr uint16_t clusterThresholdLayerOne = 2000; + constexpr uint16_t clusterThresholdOtherLayers = 4000; + + constexpr uint32_t maxNumDigis = 3 * 256 * 1024; // @PU=200 µ=530 σ=50k this is >4σ away constexpr uint16_t maxNumModules = 4000; - constexpr int32_t maxNumClustersPerModules = maxHitsInModule(); constexpr uint16_t invalidModuleId = std::numeric_limits::max() - 1; constexpr int invalidClusterId = -9999; static_assert(invalidModuleId > maxNumModules); // invalidModuleId must be > maxNumModules diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h index 0a585792ca158..89a70369fa08f 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h @@ -78,6 +78,5 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection; using TrackingRecHitSoADevicePhase2 = TrackingRecHitSoADevice; using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice; -using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice; #endif // CUDADataFormats_Track_TrackHeterogeneousT_H diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h index 6382645b7cb5b..bfac27b2b71e6 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h @@ -68,6 +68,5 @@ class TrackingRecHitSoAHost : public cms::cuda::PortableHostCollection; using TrackingRecHitSoAHostPhase2 = TrackingRecHitSoAHost; using TrackingRecHitSoAHostHIonPhase1 = TrackingRecHitSoAHost; -using TrackingRecHitSoAHostHIonPhase1 = TrackingRecHitSoAHost; #endif // CUDADataFormats_Track_TrackHeterogeneousT_H diff --git a/Configuration/PyReleaseValidation/python/relval_gpu.py b/Configuration/PyReleaseValidation/python/relval_gpu.py index 78a82711cc678..84ccdf022b908 100644 --- a/Configuration/PyReleaseValidation/python/relval_gpu.py +++ b/Configuration/PyReleaseValidation/python/relval_gpu.py @@ -67,9 +67,8 @@ # data 2023 Patatrack pixel-only triplets: RunJetMET2022D on GPU (optional) # Patatrack ECAL-only: RunJetMET2022D on GPU (optional) # Patatrack HCAL-only: RunJetMET2022D on GPU (optional) - workflows[141.008506] = ['Run3-2023_JetMET2023B_RecoPixelOnlyTripletsGPU',['RunJetMET2023B','HLTDR3_2023','RECODR3_reHLT_Patatrack_PixelOnlyTripletsGPU','HARVESTRUN3_pixelTrackingOnly']] workflows[141.008512] = ['Run3-2023_JetMET2023B_RecoECALOnlyGPU',['RunJetMET2023B','HLTDR3_2023','RECODR3_reHLT_ECALOnlyGPU','HARVESTRUN3_ECALOnly']] workflows[141.008522] = ['Run3-2023_JetMET2023B_RecoHCALOnlyGPU',['RunJetMET2023B','HLTDR3_2023','RECODR3_reHLT_HCALOnlyGPU','HARVESTRUN3_HCALOnly']] - +#2023 HIon MC Patatrack pixel-only quadruplets on HydjetQ_MinBias_5362GeV_2023_ppReco on GPU (optional) workflows[160.502] = ['',['HydjetQ_MinBias_5362GeV_2023_ppReco','DIGIHI2023PPRECO','RAWPRIMESIMHI18','RECOHI2023PPRECOMB_PatatrackGPU','MINIHI2023PROD']] \ No newline at end of file diff --git a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py index 0a00bfcd98b73..a8b952d351122 100644 --- a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py +++ b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py @@ -1490,37 +1490,6 @@ def setup_(self, step, stepName, stepDict, k, properties): offset = 0.597, ) - -class PatatrackWorkflowHI(PatatrackWorkflow): - - def condition(self, fragment, stepList, key, hasHarvest): - # select only a subset of the workflows - selected = [ - ('Hydjet_Quenched' in fragment and "PixelOnly" in self.suffix ) - ] - result = any(selected) and hasHarvest - - return result - - def setup_(self, step, stepName, stepDict, k, properties): - # skip ALCA and Nano steps (but not RecoNano or HARVESTNano for Run3) - if 'ALCA' in step or 'Nano'==step: - stepDict[stepName][k] = None - elif 'Digi' in step: - if self.__digi is None: - stepDict[stepName][k] = None - else: - stepDict[stepName][k] = merge([self.__digi, stepDict[step][k]]) - elif 'Reco' in step: - if self.__reco is None: - stepDict[stepName][k] = None - else: - stepDict[stepName][k] = merge([self.__reco, stepDict[step][k]]) - elif 'HARVEST' in step: - if self.__harvest is None: - stepDict[stepName][k] = None - else: - stepDict[stepName][k] = merge([self.__harvest, stepDict[step][k]]) # end of Patatrack workflows class UpgradeWorkflow_ProdLike(UpgradeWorkflow): @@ -2439,10 +2408,10 @@ def condition(self, fragment, stepList, key, hasHarvest): class UpgradeWorkflow_DDDDB(UpgradeWorkflow): def setup_(self, step, stepName, stepDict, k, properties): - theEra = stepDict[step][k]['--era'] - if 'Run3' in stepDict[step][k]['--era'] and '2023' not in stepDict[step][k]['--era'] and 'Fast' not in theEra and "Pb" not in theEra: + the_era = stepDict[step][k]['--era'] + if 'Run3' in the_era and '2023' not in the_era and 'Fast' not in the_era and "Pb" not in the_era: # retain any other eras - tmp_eras = stepDict[step][k]['--era'].split(',') + tmp_eras = the_era.split(',') tmp_eras[tmp_eras.index("Run3")] = 'Run3_DDD' tmp_eras = ','.join(tmp_eras) stepDict[stepName][k] = merge([{'--conditions': 'auto:phase1_2022_realistic_ddd', '--geometry': 'DB:Extended', '--era': tmp_eras}, stepDict[step][k]]) diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc index 2ab9619577039..6e2a908b59b38 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareRecHitsSoA.cc @@ -248,7 +248,6 @@ void SiPixelCompareRecHitsSoA::fillDescriptions(edm::ConfigurationDescription using SiPixelPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA; using SiPixelPhase2CompareRecHitsSoA = SiPixelCompareRecHitsSoA; using SiPixelHIonPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA; -using SiPixelHIonPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA; DEFINE_FWK_MODULE(SiPixelPhase1CompareRecHitsSoA); DEFINE_FWK_MODULE(SiPixelPhase2CompareRecHitsSoA); diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc index da7ec10f47507..03d023cf17a71 100644 --- a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc +++ b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc @@ -314,4 +314,3 @@ using SiPixelHIonPhase1CompareTrackSoA = SiPixelCompareTrackSoA void SiPixelDigisClustersFromSoAT::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; desc.add("src", edm::InputTag("siPixelDigisSoA")); - desc.add("clusterThreshold_layer1", 2000); //TODO put these somewhere - desc.add("clusterThreshold_otherLayers", 4000); + desc.add("clusterThreshold_layer1", gpuClustering::clusterThresholdLayerOne); + desc.add("clusterThreshold_otherLayers", gpuClustering::clusterThresholdOtherLayers); desc.add("produceDigis", true); desc.add("storeDigis", true); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc index 7fac305a9bbd8..82ff6aec2ea3c 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc @@ -79,10 +79,10 @@ SiPixelPhase2DigiToClusterCUDA::SiPixelPhase2DigiToClusterCUDA(const edm::Parame includeErrors_(iConfig.getParameter("IncludeErrors")), clusterThresholds_{iConfig.getParameter("clusterThreshold_layer1"), iConfig.getParameter("clusterThreshold_otherLayers"), - (float)iConfig.getParameter("ElectronPerADCGain"), - (int8_t)iConfig.getParameter("Phase2ReadoutMode"), - (uint16_t)iConfig.getParameter("Phase2DigiBaseline"), - (uint8_t)iConfig.getParameter("Phase2KinkADC")} { + static_cast(iConfig.getParameter("ElectronPerADCGain")), + static_cast(iConfig.getParameter("Phase2ReadoutMode")), + static_cast(iConfig.getParameter("Phase2DigiBaseline")), + static_cast(iConfig.getParameter("Phase2KinkADC"))} { if (includeErrors_) { digiErrorPutToken_ = produces>(); } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc index aba309f478308..8aa39335d27d3 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc @@ -6,6 +6,7 @@ // CMSSW includes #include "CUDADataFormats/Common/interface/Product.h" #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h" +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h" #include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h" @@ -98,10 +99,10 @@ SiPixelRawToClusterCUDAT::SiPixelRawToClusterCUDAT(const edm::Par useQuality_(iConfig.getParameter("UseQualityInfo")), clusterThresholds_{iConfig.getParameter("clusterThreshold_layer1"), iConfig.getParameter("clusterThreshold_otherLayers"), - (float)iConfig.getParameter("VCaltoElectronGain"), - (float)iConfig.getParameter("VCaltoElectronGain_L1"), - (float)iConfig.getParameter("VCaltoElectronOffset"), - (float)iConfig.getParameter("VCaltoElectronOffset_L1")} { + static_cast(iConfig.getParameter("VCaltoElectronGain")), + static_cast(iConfig.getParameter("VCaltoElectronGain_L1")), + static_cast(iConfig.getParameter("VCaltoElectronOffset")), + static_cast(iConfig.getParameter("VCaltoElectronOffset_L1"))} { if (includeErrors_) { digiErrorPutToken_ = produces>(); } @@ -123,8 +124,8 @@ void SiPixelRawToClusterCUDAT::fillDescriptions(edm::Configuratio desc.addOptionalNode(edm::ParameterDescription("MaxFEDWords", 0, true), false) ->setComment("This parameter is obsolete and will be ignored."); //Clustering Thresholds - desc.add("clusterThreshold_layer1", 2000); - desc.add("clusterThreshold_otherLayers", 4000); + desc.add("clusterThreshold_layer1", gpuClustering::clusterThresholdLayerOne); + desc.add("clusterThreshold_otherLayers", gpuClustering::clusterThresholdOtherLayers); desc.add("VCaltoElectronGain", 47.f); desc.add("VCaltoElectronGain_L1", 50.f); desc.add("VCaltoElectronOffset", -60.f); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h index cfd6efb3eef2d..a96cd0bcc5c15 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -27,11 +27,6 @@ namespace gpuClustering { uint32_t numElements) { constexpr int32_t maxNumClustersPerModules = TrackerTraits::maxNumClustersPerModules; - static_assert( - maxNumClustersPerModules <= 2048, - "\nclusterChargeCut is limited to 2048 clusters per module. \nHere maxNumClustersPerModules is set to be %d. " - "\nIf you need maxNumClustersPerModules to be higher \nyou will need to fix the blockPrefixScans."); - __shared__ int32_t charge[maxNumClustersPerModules]; __shared__ uint8_t ok[maxNumClustersPerModules]; __shared__ uint16_t newclusId[maxNumClustersPerModules]; @@ -132,13 +127,15 @@ namespace gpuClustering { cms::cuda::blockPrefixScan(newclusId, newclusId, minClust, ws); if constexpr (maxNumClustersPerModules > maxThreads) //only if needed { - //TODO: most probably there's a smarter implementation for this - if (nclus > maxThreads) { - cms::cuda::blockPrefixScan(newclusId + maxThreads, newclusId + maxThreads, nclus - maxThreads, ws); - for (auto i = threadIdx.x + maxThreads; i < nclus; i += blockDim.x) { - int prevBlockEnd = ((i / maxThreads) * maxThreads) - 1; + for (uint32_t offset = maxThreads; offset < nclus; offset += maxThreads) { + cms::cuda::blockPrefixScan(newclusId + offset, newclusId + offset, nclus - offset, ws); + + for (uint32_t i = threadIdx.x + offset; i < nclus; i += blockDim.x) { + uint32_t prevBlockEnd = ((i / maxThreads) * maxThreads) - 1; newclusId[i] += newclusId[prevBlockEnd]; } + + __syncthreads(); } } assert(nclus > newclusId[nclus - 1]); diff --git a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py index 2f193f7c861f4..263e0f77e5fbf 100644 --- a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py +++ b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py @@ -19,10 +19,12 @@ from RecoLocalTracker.SiPixelClusterizer.siPixelRawToClusterCUDAHIonPhase1_cfi import siPixelRawToClusterCUDAHIonPhase1 as _siPixelRawToClusterCUDAHIonPhase1 siPixelClustersPreSplittingCUDA = _siPixelRawToClusterCUDA.clone() -# HIon modifiers +# HIon Modifiers from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA +# Phase 2 Tracker Modifier +from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker -pp_on_AA.toReplaceWith(siPixelClustersPreSplittingCUDA, _siPixelRawToClusterCUDAHIonPhase1.clone()) +(pp_on_AA & ~phase2_tracker).toReplaceWith(siPixelClustersPreSplittingCUDA, _siPixelRawToClusterCUDAHIonPhase1.clone()) run3_common.toModify(siPixelClustersPreSplittingCUDA, # use the pixel channel calibrations scheme for Run 3 @@ -40,14 +42,12 @@ siPixelDigisClustersPreSplitting = _siPixelDigisClustersFromSoAPhase1.clone() from RecoLocalTracker.SiPixelClusterizer.siPixelDigisClustersFromSoAHIonPhase1_cfi import siPixelDigisClustersFromSoAHIonPhase1 as _siPixelDigisClustersFromSoAHIonPhase1 -pp_on_AA.toReplaceWith(siPixelDigisClustersPreSplitting, _siPixelDigisClustersFromSoAHIonPhase1.clone()) +(pp_on_AA & ~phase2_tracker).toReplaceWith(siPixelDigisClustersPreSplitting, _siPixelDigisClustersFromSoAHIonPhase1.clone()) run3_common.toModify(siPixelDigisClustersPreSplitting, clusterThreshold_layer1 = 4000) -from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker - gpu.toReplaceWith(siPixelClustersPreSplittingTask, cms.Task( # conditions used *only* by the modules running on GPU siPixelROCsStatusAndMappingWrapperESProducer, diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h index a8f7dae43b5d7..c0291ed9f32f8 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h +++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h @@ -19,6 +19,7 @@ #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterThresholds.h" +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "Geometry/CommonTopologies/interface/SimplePixelTopology.h" @@ -31,7 +32,8 @@ int main(void) { using pixelTopology::Phase1; constexpr int numElements = 256 * maxNumModules; - const SiPixelClusterThresholds clusterThresholds(2000, 4000, 0.f, 0.f, 0.f, 0.f); + const SiPixelClusterThresholds clusterThresholds( + clusterThresholdLayerOne, clusterThresholdOtherLayers, 0.f, 0.f, 0.f, 0.f); // these in reality are already on GPU auto h_raw = std::make_unique(numElements); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 1da748d8dcd8b..8dc6ae93018ea 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -117,7 +117,7 @@ void SiPixelRecHitSoAFromLegacyT::produce(edm::StreamID streamID, std::vector, SiPixelCluster>> clusterRef; - constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule(); + constexpr uint32_t maxHitsInModule = TrackerTraits::maxNumClustersPerModules; cms::cuda::PortableHostCollection> clusters_h(nModules + 1); diff --git a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py index c4ee471950b93..f45b41861995d 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py +++ b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py @@ -27,7 +27,7 @@ from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromLegacyHIonPhase1_cfi import siPixelRecHitSoAFromLegacyHIonPhase1 as _siPixelRecHitsPreSplittingSoAHIonPhase1 siPixelRecHitsPreSplittingCPU = _siPixelRecHitsPreSplittingSoA.clone(convertToLegacy=True) -pp_on_AA.toReplaceWith(siPixelRecHitsPreSplittingCPU, _siPixelRecHitsPreSplittingSoAHIonPhase1.clone(convertToLegacy=True, CPE = cms.string('PixelCPEFastHIonPhase1'))) +(pp_on_AA & ~phase2_tracker).toReplaceWith(siPixelRecHitsPreSplittingCPU, _siPixelRecHitsPreSplittingSoAHIonPhase1.clone(convertToLegacy=True, CPE = cms.string('PixelCPEFastHIonPhase1'))) phase2_tracker.toReplaceWith(siPixelRecHitsPreSplittingCPU, _siPixelRecHitsPreSplittingSoAPhase2.clone(convertToLegacy=True, CPE = cms.string('PixelCPEFastPhase2'))) # modifier used to prompt patatrack pixel tracks reconstruction on cpu @@ -47,7 +47,6 @@ ) # reconstruct the pixel rechits on the gpu - from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDAPhase1_cfi import siPixelRecHitCUDAPhase1 as _siPixelRecHitCUDAPhase1 from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDAPhase2_cfi import siPixelRecHitCUDAPhase2 as _siPixelRecHitCUDAPhase2 from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDAHIonPhase1_cfi import siPixelRecHitCUDAHIonPhase1 as _siPixelRecHitCUDAHIonPhase1 @@ -55,7 +54,7 @@ siPixelRecHitsPreSplittingCUDA = _siPixelRecHitCUDAPhase1.clone( beamSpot = "offlineBeamSpotToCUDA" ) -pp_on_AA.toReplaceWith(siPixelRecHitsPreSplittingCUDA,_siPixelRecHitCUDAHIonPhase1.clone( +(pp_on_AA & ~phase2_tracker).toReplaceWith(siPixelRecHitsPreSplittingCUDA,_siPixelRecHitCUDAHIonPhase1.clone( beamSpot = "offlineBeamSpotToCUDA" )) phase2_tracker.toReplaceWith(siPixelRecHitsPreSplittingCUDA,_siPixelRecHitCUDAPhase2.clone( @@ -71,15 +70,15 @@ )), ) -pp_on_AA.toModify(siPixelRecHitsPreSplittingSoA, -cpu = cms.EDAlias( +(pp_on_AA & ~phase2_tracker).toModify(siPixelRecHitsPreSplittingSoA, + cpu = cms.EDAlias( siPixelRecHitsPreSplittingCPU = cms.VPSet( cms.PSet(type = cms.string("pixelTopologyHIonPhase1TrackingRecHitSoAHost")), cms.PSet(type = cms.string("uintAsHostProduct")) ))) phase2_tracker.toModify(siPixelRecHitsPreSplittingSoA, -cpu = cms.EDAlias( + cpu = cms.EDAlias( siPixelRecHitsPreSplittingCPU = cms.VPSet( cms.PSet(type = cms.string("pixelTopologyPhase2TrackingRecHitSoAHost")), cms.PSet(type = cms.string("uintAsHostProduct")) diff --git a/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py b/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py index 0a3e57a9ff96a..c08a0987d3f59 100644 --- a/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py +++ b/RecoTracker/Configuration/python/RecoPixelVertexing_cff.py @@ -34,7 +34,7 @@ PtMin = 2.0 )) -pp_on_AA.toModify(pixelVerticesSoA,cpu = _pixelVerticesCUDAHIonPhase1.clone( +(pp_on_AA & ~phase2_tracker).toModify(pixelVerticesSoA,cpu = _pixelVerticesCUDAHIonPhase1.clone( pixelTrackSrc = "pixelTracksSoA", doSplitting = False, onGPU = False, @@ -70,7 +70,7 @@ PtMin = 2.0 )) -pp_on_AA.toReplaceWith(pixelVerticesCUDA,_pixelVerticesCUDAHIonPhase1.clone( +(pp_on_AA & ~phase2_tracker).toReplaceWith(pixelVerticesCUDA,_pixelVerticesCUDAHIonPhase1.clone( pixelTrackSrc = "pixelTracksCUDA", doSplitting = False, onGPU = True diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc index 3c4125a75b6b1..9ab908a037bd7 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cc @@ -73,7 +73,7 @@ void CAHitNtupletGeneratorKernelsCPU::buildDoublets(const HitsCon this->isOuterHitOfCell_, nActualPairs, this->params_.caParams_.maxNumberOfDoublets_, - this->cellCuts_.get()); + this->device_cellCuts_.get()); } template diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu index b13cdcda4d28d..efb2a2e17715c 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu @@ -236,7 +236,7 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(const HitsCon this->isOuterHitOfCell_, nActualPairs, this->params_.caParams_.maxNumberOfDoublets_, - this->cellCuts_.get()); + this->device_cellCuts_.get()); cudaCheck(cudaGetLastError()); #ifdef GPU_DEBUG diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h index 55dd7412a46bb..0865fa5cbc46a 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h @@ -247,7 +247,7 @@ class CAHitNtupletGeneratorKernels { protected: Counters* counters_ = nullptr; - // CellCuts* cellCuts_ = nullptr; + // workspace unique_ptr cellStorage_; unique_ptr device_theCellNeighbors_; @@ -264,7 +264,7 @@ class CAHitNtupletGeneratorKernels { unique_ptr device_hitToTupleStorage_; typename HitToTuple::View hitToTupleView_; - unique_ptr cellCuts_; + unique_ptr device_cellCuts_; cms::cuda::AtomicPairCounter* device_hitToTuple_apc_ = nullptr; diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc index 01cca3187d5f4..6acff4abbd531 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc @@ -41,15 +41,15 @@ void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU(int32_t nHits this->device_hitToTuple_apc_ = (cms::cuda::AtomicPairCounter*)this->device_storage_.get() + 1; this->device_nCells_ = (uint32_t*)(this->device_storage_.get() + 2); - this->cellCuts_ = Traits::template make_unique(stream); + this->device_cellCuts_ = Traits::template make_unique(stream); // FIXME: consider collapsing these 3 in one adhoc kernel if constexpr (std::is_same::value) { cudaCheck(cudaMemsetAsync(this->device_nCells_, 0, sizeof(uint32_t), stream)); cudaCheck(cudaMemcpyAsync( - this->cellCuts_.get(), &(this->params_.cellCuts_), sizeof(CellCuts), cudaMemcpyDefault, stream)); + this->device_cellCuts_.get(), &(this->params_.cellCuts_), sizeof(CellCuts), cudaMemcpyDefault, stream)); } else { *(this->device_nCells_) = 0; - *(this->cellCuts_.get()) = this->params_.cellCuts_; + *(this->device_cellCuts_.get()) = this->params_.cellCuts_; } cms::cuda::launchZero(this->device_tupleMultiplicity_.get(), stream); cms::cuda::launchZero(this->hitToTupleView_, stream); // we may wish to keep it in the edm diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc index 5e305bebd0e95..cf5593eed6543 100644 --- a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc +++ b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc @@ -64,12 +64,12 @@ namespace { return CAParamsT{{ cfg.getParameter("maxNumberOfDoublets"), cfg.getParameter("minHitsPerNtuplet"), - (float)cfg.getParameter("ptmin"), - (float)cfg.getParameter("CAThetaCutBarrel"), - (float)cfg.getParameter("CAThetaCutForward"), - (float)cfg.getParameter("hardCurvCut"), - (float)cfg.getParameter("dcaCutInnerTriplet"), - (float)cfg.getParameter("dcaCutOuterTriplet"), + static_cast(cfg.getParameter("ptmin")), + static_cast(cfg.getParameter("CAThetaCutBarrel")), + static_cast(cfg.getParameter("CAThetaCutForward")), + static_cast(cfg.getParameter("hardCurvCut")), + static_cast(cfg.getParameter("dcaCutInnerTriplet")), + static_cast(cfg.getParameter("dcaCutOuterTriplet")), }}; }; @@ -100,21 +100,21 @@ namespace { static constexpr CAParamsT makeCACuts(edm::ParameterSet const& cfg) { return CAParamsT{{cfg.getParameter("maxNumberOfDoublets"), cfg.getParameter("minHitsPerNtuplet"), - (float)cfg.getParameter("ptmin"), - (float)cfg.getParameter("CAThetaCutBarrel"), - (float)cfg.getParameter("CAThetaCutForward"), - (float)cfg.getParameter("hardCurvCut"), - (float)cfg.getParameter("dcaCutInnerTriplet"), - (float)cfg.getParameter("dcaCutOuterTriplet")}, + static_cast(cfg.getParameter("ptmin")), + static_cast(cfg.getParameter("CAThetaCutBarrel")), + static_cast(cfg.getParameter("CAThetaCutForward")), + static_cast(cfg.getParameter("hardCurvCut")), + static_cast(cfg.getParameter("dcaCutInnerTriplet")), + static_cast(cfg.getParameter("dcaCutOuterTriplet"))}, {(bool)cfg.getParameter("includeFarForwards")}}; } static constexpr pixelTrack::QualityCutsT makeQualityCuts(edm::ParameterSet const& pset) { return pixelTrack::QualityCutsT{ - (float)pset.getParameter("maxChi2"), - (float)pset.getParameter("minPt"), - (float)pset.getParameter("maxTip"), - (float)pset.getParameter("maxZip"), + static_cast(pset.getParameter("maxChi2")), + static_cast(pset.getParameter("minPt")), + static_cast(pset.getParameter("maxTip")), + static_cast(pset.getParameter("maxZip")), }; } }; @@ -197,8 +197,8 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::Pa trackQualityCuts.add("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm"); trackQualityCuts.add("quadrupletMaxZip", 12.)->setComment("Max |Zip| for quadruplets, in cm"); - desc.add>("phiCuts", - std::vector(phase1PixelTopology::phicuts, std::end(phase1PixelTopology::phicuts))) + desc.add>( + "phiCuts", std::vector(std::begin(phase1PixelTopology::phicuts), std::end(phase1PixelTopology::phicuts))) ->setComment("Cuts in phi for cells"); desc.add("trackQualityCuts", trackQualityCuts) @@ -230,8 +230,8 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm trackQualityCuts.add("quadrupletMaxTip", 0.5)->setComment("Max |Tip| for quadruplets, in cm"); trackQualityCuts.add("quadrupletMaxZip", 6.)->setComment("Max |Zip| for quadruplets, in cm"); - desc.add>("phiCuts", - std::vector(phase1PixelTopology::phicuts, std::end(phase1PixelTopology::phicuts))) + desc.add>( + "phiCuts", std::vector(std::begin(phase1PixelTopology::phicuts), std::end(phase1PixelTopology::phicuts))) ->setComment("Cuts in phi for cells"); desc.add("trackQualityCuts", trackQualityCuts) @@ -256,8 +256,8 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::Pa trackQualityCuts.add("maxTip", 0.3)->setComment("Max |Tip| in cm"); trackQualityCuts.add("maxZip", 12.)->setComment("Max |Zip|, in cm"); - desc.add>("phiCuts", - std::vector(phase2PixelTopology::phicuts, std::end(phase2PixelTopology::phicuts))) + desc.add>( + "phiCuts", std::vector(std::begin(phase2PixelTopology::phicuts), std::end(phase2PixelTopology::phicuts))) ->setComment("Cuts in phi for cells"); desc.add("trackQualityCuts", trackQualityCuts) diff --git a/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h b/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h index 7e97bc29b04d1..472fa7117fd77 100644 --- a/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h +++ b/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h @@ -50,7 +50,7 @@ namespace gpuPixelDoublets { CellCutsT* const cuts) { doubletsFromHisto( - nActualPairs, maxNumOfDoublets, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, cuts); + nActualPairs, maxNumOfDoublets, cells, nCells, cellNeighbors, cellTracks, hh, isOuterHitOfCell, *cuts); } } // namespace gpuPixelDoublets diff --git a/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h b/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h index c6f6dcfa0f822..b86ba09949416 100644 --- a/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h +++ b/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h @@ -40,6 +40,7 @@ namespace gpuPixelDoublets { using T = TrackerTraits; CellCutsT() = default; + CellCutsT(const bool doClusterCut, const bool doZ0Cut, const bool doPtCut, @@ -126,16 +127,16 @@ namespace gpuPixelDoublets { CellTracksVector* cellTracks, HitsConstView hh, OuterHitOfCell isOuterHitOfCell, - CellCutsT const* cuts) { + CellCutsT const& cuts) { // ysize cuts (z in the barrel) times 8 // these are used if doClusterCut is true - const bool doClusterCut = cuts->doClusterCut_; - const bool doZ0Cut = cuts->doZ0Cut_; - const bool doPtCut = cuts->doPtCut_; + const bool doClusterCut = cuts.doClusterCut_; + const bool doZ0Cut = cuts.doZ0Cut_; + const bool doPtCut = cuts.doPtCut_; - const float z0cut = cuts->z0Cut_; // cm - const float hardPtCut = cuts->ptCut_; // GeV + const float z0cut = cuts.z0Cut_; // cm + const float hardPtCut = cuts.ptCut_; // GeV // cm (1 GeV track has 1 GeV/c / (e * 3.8T) ~ 87 cm radius in a 3.8T field) const float minRadius = hardPtCut * 87.78f; const float minRadius2T4 = 4.f * minRadius * minRadius; @@ -206,7 +207,7 @@ namespace gpuPixelDoublets { if (mez < TrackerTraits::minz[pairLayerId] || mez > TrackerTraits::maxz[pairLayerId]) continue; - if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts->clusterCut(hh, i)) + if (doClusterCut && outer > pixelTopology::last_barrel_layer && cuts.clusterCut(hh, i)) continue; auto mep = hh[i].iphi(); @@ -226,7 +227,7 @@ namespace gpuPixelDoublets { return dr > TrackerTraits::maxr[pairLayerId] || dr < 0 || std::abs((mez * ro - mer * zo)) > z0cut * dr; }; - auto iphicut = cuts->phiCuts[pairLayerId]; + auto iphicut = cuts.phiCuts[pairLayerId]; auto kl = PhiBinner::bin(int16_t(mep - iphicut)); auto kh = PhiBinner::bin(int16_t(mep + iphicut)); @@ -264,7 +265,7 @@ namespace gpuPixelDoublets { if (idphi > iphicut) continue; - if (doClusterCut && cuts->zSizeCut(hh, i, oi)) + if (doClusterCut && cuts.zSizeCut(hh, i, oi)) continue; if (doPtCut && ptcut(oi, idphi)) continue; diff --git a/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py b/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py index 35051e7b50cf6..91eb380a33da9 100644 --- a/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py +++ b/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py @@ -175,7 +175,7 @@ onGPU = False )) -pp_on_AA.toModify(pixelTracksSoA,cpu = _pixelTracksCUDAHIonPhase1.clone( +(pp_on_AA & ~phase2_tracker).toModify(pixelTracksSoA,cpu = _pixelTracksCUDAHIonPhase1.clone( pixelRecHitSrc = "siPixelRecHitsPreSplittingSoA", onGPU = False )) @@ -185,7 +185,7 @@ onGPU = True, )) -pp_on_AA.toReplaceWith(pixelTracksCUDA,_pixelTracksCUDAHIonPhase1.clone( +(pp_on_AA & ~phase2_tracker).toReplaceWith(pixelTracksCUDA,_pixelTracksCUDAHIonPhase1.clone( pixelRecHitSrc = "siPixelRecHitsPreSplittingCUDA", onGPU = True, )) diff --git a/Validation/RecoTrack/python/PostProcessorTracker_cfi.py b/Validation/RecoTrack/python/PostProcessorTracker_cfi.py index ac9686b547701..c1ae8bb670d58 100644 --- a/Validation/RecoTrack/python/PostProcessorTracker_cfi.py +++ b/Validation/RecoTrack/python/PostProcessorTracker_cfi.py @@ -352,17 +352,25 @@ def _addNoFlow(module): from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA _defaultSubdirsHIon = _defaultSubdirs + ["Tracking/HIPixelTrack/*"] -_defaultSubdirsSummaryHIon = [e.replace("/*","") for e in _defaultSubdirsHIon] -pp_on_AA.toModify(postProcessorTrack,subDirs = _defaultSubdirsHIon) -pp_on_AA.toModify(postProcessorTrack2D,subDirs = _defaultSubdirsHIon) -pp_on_AA.toModify(postProcessorTrackSummary,subDirs = _defaultSubdirsHIon) +(pp_on_AA & ~phase2_tracker).toModify(postProcessorTrack,subDirs = _defaultSubdirsHIon) +(pp_on_AA & ~phase2_tracker).toModify(postProcessorTrack2D,subDirs = _defaultSubdirsHIon) +(pp_on_AA & ~phase2_tracker).toModify(postProcessorTrackSummary,subDirs = _defaultSubdirsHIon) postProcessorTrackTrackingOnly = postProcessorTrack.clone() -postProcessorTrackTrackingOnly.subDirs.extend(["Tracking/TrackBHadron/*", "Tracking/TrackSeeding/*", "Tracking/HIPixelTrack/*", "Tracking/PixelTrack/*", "Tracking/PixelTrackFromPV/*", "Tracking/PixelTrackFromPVAllTP/*", "Tracking/PixelTrackBHadron/*"]) +postProcessorTrackTrackingOnly.subDirs.extend(["Tracking/TrackBHadron/*", "Tracking/TrackSeeding/*", "Tracking/PixelTrack/*", "Tracking/PixelTrackFromPV/*", "Tracking/PixelTrackFromPVAllTP/*", "Tracking/PixelTrackBHadron/*"]) postProcessorTrackSummaryTrackingOnly = postProcessorTrackSummary.clone() -postProcessorTrackSummaryTrackingOnly.subDirs.extend(["Tracking/TrackBHadron", "Tracking/TrackSeeding", "Tracking/HIPixelTrack", "Tracking/PixelTrack", "Tracking/PixelTrackFromPV", "Tracking/PixelTrackFromPVAllTP", "Tracking/PixelTrackBHadron"]) +postProcessorTrackSummaryTrackingOnly.subDirs.extend(["Tracking/TrackBHadron", "Tracking/TrackSeeding", "Tracking/PixelTrack", "Tracking/PixelTrackFromPV", "Tracking/PixelTrackFromPVAllTP", "Tracking/PixelTrackBHadron"]) + +postProcessorTrackTrackingOnlyHIon = postProcessorTrackTrackingOnly.clone() +postProcessorTrackTrackingOnlyHIon.subDirs.extend(["Tracking/HIPixelTrack/*"]) + +postProcessorTrackSummaryTrackingOnlyHIon = postProcessorTrackSummaryTrackingOnly.clone() +postProcessorTrackSummaryTrackingOnlyHIon.subDirs.extend(["Tracking/HIPixelTrack"]) + +(pp_on_AA & ~phase2_tracker).toReplaceWith(postProcessorTrackTrackingOnly,postProcessorTrackTrackingOnlyHIon) +(pp_on_AA & ~phase2_tracker).toReplaceWith(postProcessorTrackSummaryTrackingOnly,postProcessorTrackSummaryTrackingOnlyHIon) postProcessorTrackSequenceTrackingOnly = cms.Sequence( postProcessorTrackTrackingOnly+