Addressing comments

AdrianoDee · AdrianoDee · commit 0c04805cf879 · 2023-07-11T16:29:54.000+02:00
diff --git 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<uint16_t>::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
@@ -78,6 +78,5 @@ class TrackingRecHitSoADevice : public cms::cuda::PortableDeviceCollection<Track
 using TrackingRecHitSoADevicePhase1 = TrackingRecHitSoADevice<pixelTopology::Phase1>;
 using TrackingRecHitSoADevicePhase2 = TrackingRecHitSoADevice<pixelTopology::Phase2>;
 using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice<pixelTopology::HIonPhase1>;
-using TrackingRecHitSoADeviceHIonPhase1 = TrackingRecHitSoADevice<pixelTopology::HIonPhase1>;
 
 #endif  // CUDADataFormats_Track_TrackHeterogeneousT_H
diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h
@@ -68,6 +68,5 @@ class TrackingRecHitSoAHost : public cms::cuda::PortableHostCollection<TrackingR
 using TrackingRecHitSoAHostPhase1 = TrackingRecHitSoAHost<pixelTopology::Phase1>;
 using TrackingRecHitSoAHostPhase2 = TrackingRecHitSoAHost<pixelTopology::Phase2>;
 using TrackingRecHitSoAHostHIonPhase1 = TrackingRecHitSoAHost<pixelTopology::HIonPhase1>;
-using TrackingRecHitSoAHostHIonPhase1 = TrackingRecHitSoAHost<pixelTopology::HIonPhase1>;
 
 #endif  // CUDADataFormats_Track_TrackHeterogeneousT_H
diff --git 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']]
diff --git 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
@@ -248,7 +248,6 @@ void SiPixelCompareRecHitsSoA<T>::fillDescriptions(edm::ConfigurationDescription
 using SiPixelPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA<pixelTopology::Phase1>;
 using SiPixelPhase2CompareRecHitsSoA = SiPixelCompareRecHitsSoA<pixelTopology::Phase2>;
 using SiPixelHIonPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA<pixelTopology::HIonPhase1>;
-using SiPixelHIonPhase1CompareRecHitsSoA = SiPixelCompareRecHitsSoA<pixelTopology::HIonPhase1>;
 
 DEFINE_FWK_MODULE(SiPixelPhase1CompareRecHitsSoA);
 DEFINE_FWK_MODULE(SiPixelPhase2CompareRecHitsSoA);
diff --git a/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc b/DQM/SiPixelHeterogeneous/plugins/SiPixelCompareTrackSoA.cc
@@ -314,4 +314,3 @@ using SiPixelHIonPhase1CompareTrackSoA = SiPixelCompareTrackSoA<pixelTopology::H
 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
@@ -209,4 +209,3 @@ using SiPixelHIonPhase1MonitorRecHitsSoA = SiPixelMonitorRecHitsSoA<pixelTopolog
 DEFINE_FWK_MODULE(SiPixelPhase1MonitorRecHitsSoA);
 DEFINE_FWK_MODULE(SiPixelPhase2MonitorRecHitsSoA);
 DEFINE_FWK_MODULE(SiPixelHIonPhase1MonitorRecHitsSoA);
-
diff --git a/DQM/SiPixelHeterogeneous/python/SiPixelHeterogenousDQM_FirstStep_cff.py b/DQM/SiPixelHeterogeneous/python/SiPixelHeterogenousDQM_FirstStep_cff.py
@@ -20,7 +20,7 @@
 from Configuration.ProcessModifiers.pp_on_AA_cff import pp_on_AA
 
 _monitorpixelSoARecHitsSourceHIon = cms.Sequence(siPixelHIonPhase1MonitorRecHitsSoA * siPixelHIonPhase1MonitorTrackSoA * siPixelMonitorVertexSoA)
-pp_on_AA.toReplaceWith(monitorpixelSoASource, _monitorpixelSoARecHitsSourceHIon)
+(pp_on_AA & ~phase2_tracker).toReplaceWith(monitorpixelSoASource, _monitorpixelSoARecHitsSourceHIon)
 
 #Define the sequence for GPU vs CPU validation
 #This should run:- individual monitor for the 2 collections + comparison module
diff --git a/Geometry/CommonTopologies/interface/SimplePixelTopology.h b/Geometry/CommonTopologies/interface/SimplePixelTopology.h
@@ -287,6 +287,8 @@ namespace phase2PixelTopology {
 }  // namespace phase2PixelTopology
 
 namespace phase1HIonPixelTopology {
+  // Storing here the needed constants different w.r.t. pp Phase1 topology.
+  // All the other defined by inheritance in the HIon topology struct.
   using pixelTopology::phi0p09;
 
   constexpr uint32_t maxNumClustersPerModules = 2048;
@@ -540,6 +542,9 @@ namespace pixelTopology {
   };
 
   struct HIonPhase1 : public Phase1 {
+    // Storing here the needed constants different w.r.t. pp Phase1 topology.
+    // All the other defined by inheritance in the HIon topology struct.
+
     using tindex_type = uint32_t;  // for tuples
 
     static constexpr uint32_t maxCellNeighbors = 90;
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc
@@ -62,8 +62,8 @@ template <typename TrackerTraits>
 void SiPixelDigisClustersFromSoAT<TrackerTraits>::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
   edm::ParameterSetDescription desc;
   desc.add<edm::InputTag>("src", edm::InputTag("siPixelDigisSoA"));
-  desc.add<int>("clusterThreshold_layer1", 2000);  //TODO put these somewhere
-  desc.add<int>("clusterThreshold_otherLayers", 4000);
+  desc.add<int>("clusterThreshold_layer1", gpuClustering::clusterThresholdLayerOne);
+  desc.add<int>("clusterThreshold_otherLayers", gpuClustering::clusterThresholdOtherLayers);
   desc.add<bool>("produceDigis", true);
   desc.add<bool>("storeDigis", true);
 
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelPhase2DigiToClusterCUDA.cc
@@ -79,10 +79,10 @@ SiPixelPhase2DigiToClusterCUDA::SiPixelPhase2DigiToClusterCUDA(const edm::Parame
       includeErrors_(iConfig.getParameter<bool>("IncludeErrors")),
       clusterThresholds_{iConfig.getParameter<int32_t>("clusterThreshold_layer1"),
                          iConfig.getParameter<int32_t>("clusterThreshold_otherLayers"),
-                         (float)iConfig.getParameter<double>("ElectronPerADCGain"),
-                         (int8_t)iConfig.getParameter<int>("Phase2ReadoutMode"),
-                         (uint16_t)iConfig.getParameter<uint32_t>("Phase2DigiBaseline"),
-                         (uint8_t)iConfig.getParameter<uint32_t>("Phase2KinkADC")} {
+                         static_cast<float>(iConfig.getParameter<double>("ElectronPerADCGain")),
+                         static_cast<int8_t>(iConfig.getParameter<int>("Phase2ReadoutMode")),
+                         static_cast<uint16_t>(iConfig.getParameter<uint32_t>("Phase2DigiBaseline")),
+                         static_cast<uint8_t>(iConfig.getParameter<uint32_t>("Phase2KinkADC"))} {
   if (includeErrors_) {
     digiErrorPutToken_ = produces<cms::cuda::Product<SiPixelDigiErrorsCUDA>>();
   }
diff --git 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<TrackerTraits>::SiPixelRawToClusterCUDAT(const edm::Par
       useQuality_(iConfig.getParameter<bool>("UseQualityInfo")),
       clusterThresholds_{iConfig.getParameter<int32_t>("clusterThreshold_layer1"),
                          iConfig.getParameter<int32_t>("clusterThreshold_otherLayers"),
-                         (float)iConfig.getParameter<double>("VCaltoElectronGain"),
-                         (float)iConfig.getParameter<double>("VCaltoElectronGain_L1"),
-                         (float)iConfig.getParameter<double>("VCaltoElectronOffset"),
-                         (float)iConfig.getParameter<double>("VCaltoElectronOffset_L1")} {
+                         static_cast<float>(iConfig.getParameter<double>("VCaltoElectronGain")),
+                         static_cast<float>(iConfig.getParameter<double>("VCaltoElectronGain_L1")),
+                         static_cast<float>(iConfig.getParameter<double>("VCaltoElectronOffset")),
+                         static_cast<float>(iConfig.getParameter<double>("VCaltoElectronOffset_L1"))} {
   if (includeErrors_) {
     digiErrorPutToken_ = produces<cms::cuda::Product<SiPixelDigiErrorsCUDA>>();
   }
@@ -123,8 +124,8 @@ void SiPixelRawToClusterCUDAT<TrackerTraits>::fillDescriptions(edm::Configuratio
   desc.addOptionalNode(edm::ParameterDescription<uint32_t>("MaxFEDWords", 0, true), false)
       ->setComment("This parameter is obsolete and will be ignored.");
   //Clustering Thresholds
-  desc.add<int32_t>("clusterThreshold_layer1", 2000);
-  desc.add<int32_t>("clusterThreshold_otherLayers", 4000);
+  desc.add<int32_t>("clusterThreshold_layer1", gpuClustering::clusterThresholdLayerOne);
+  desc.add<int32_t>("clusterThreshold_otherLayers", gpuClustering::clusterThresholdOtherLayers);
   desc.add<double>("VCaltoElectronGain", 47.f);
   desc.add<double>("VCaltoElectronGain_L1", 50.f);
   desc.add<double>("VCaltoElectronOffset", -60.f);
diff --git 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
@@ -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
@@ -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<uint32_t[]>(numElements);
diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc
@@ -117,7 +117,7 @@ void SiPixelRecHitSoAFromLegacyT<TrackerTraits>::produce(edm::StreamID streamID,
 
   std::vector<edm::Ref<edmNew::DetSetVector<SiPixelCluster>, SiPixelCluster>> clusterRef;
 
-  constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule();
+  constexpr uint32_t maxHitsInModule = TrackerTraits::maxNumClustersPerModules;
 
   cms::cuda::PortableHostCollection<SiPixelClustersCUDALayout<>> clusters_h(nModules + 1);
 
diff --git 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,15 +47,14 @@
 )
 
 # 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(
+(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
@@ -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
@@ -73,7 +73,7 @@ void CAHitNtupletGeneratorKernelsCPU<TrackerTraits>::buildDoublets(const HitsCon
                                       this->isOuterHitOfCell_,
                                       nActualPairs,
                                       this->params_.caParams_.maxNumberOfDoublets_,
-                                      this->cellCuts_.get());
+                                      this->device_cellCuts_.get());
 }
 
 template <typename TrackerTraits>
diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.cu
diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernels.h
diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorKernelsAlloc.cc
diff --git a/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc b/RecoTracker/PixelSeeding/plugins/CAHitNtupletGeneratorOnGPU.cc
diff --git a/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h b/RecoTracker/PixelSeeding/plugins/gpuPixelDoublets.h
diff --git a/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h b/RecoTracker/PixelSeeding/plugins/gpuPixelDoubletsAlgos.h
diff --git a/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py b/RecoTracker/PixelTrackFitting/python/PixelTracks_cff.py
diff --git a/Validation/RecoTrack/python/PostProcessorTracker_cfi.py b/Validation/RecoTrack/python/PostProcessorTracker_cfi.py
