Synchronise with CMSSW_10_2_6

CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h

@@ -0,0 +1,14 @@
+#ifndef CalibTracker_Records_SiPixelGainCalibrationForHLTGPURcd_h
+#define CalibTracker_Records_SiPixelGainCalibrationForHLTGPURcd_h
+
+#include "FWCore/Framework/interface/EventSetupRecordImplementation.h"
+#include "FWCore/Framework/interface/DependentRecordImplementation.h"
+
+#include "CondFormats/DataRecord/interface/SiPixelGainCalibrationForHLTRcd.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+
+#include "boost/mpl/vector.hpp"
+
+class SiPixelGainCalibrationForHLTGPURcd : public edm::eventsetup::DependentRecordImplementation<SiPixelGainCalibrationForHLTGPURcd, boost::mpl::vector<SiPixelGainCalibrationForHLTRcd, TrackerDigiGeometryRecord> > {};
+
+#endif

CalibTracker/Records/src/SiPixelGainCalibrationForHLTGPURcd.cc

@@ -0,0 +1,5 @@
+#include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+EVENTSETUP_RECORD_REG(SiPixelGainCalibrationForHLTGPURcd);

CalibTracker/SiPixelESProducers/BuildFile.xml

@@ -7,7 +7,9 @@
 <use   name="DataFormats/SiPixelDigi"/>
 <use   name="CalibTracker/Records"/>
 <use   name="MagneticField/VolumeBasedEngine"/>
+<use   name="HeterogeneousCore/CUDACore"/>
 <use   name="boost"/>
+<use   name="cuda-api-wrappers"/>
 <export>
   <lib   name="1"/>
 </export>

CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h

@@ -0,0 +1,32 @@
+#ifndef CalibTracker_SiPixelESProducers_SiPixelGainCalibrationForHLTGPU_H
+#define CalibTracker_SiPixelESProducers_SiPixelGainCalibrationForHLTGPU_H
+
+#include "HeterogeneousCore/CUDACore/interface/CUDAESProduct.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+
+#include <cuda/api_wrappers.h>
+
+class SiPixelGainCalibrationForHLT;
+class SiPixelGainForHLTonGPU;
+struct SiPixelGainForHLTonGPU_DecodingStructure;
+class TrackerGeometry;
+
+class SiPixelGainCalibrationForHLTGPU {
+public:
+  explicit SiPixelGainCalibrationForHLTGPU(const SiPixelGainCalibrationForHLT& gains, const TrackerGeometry& geom);
+  ~SiPixelGainCalibrationForHLTGPU();
+
+  const SiPixelGainForHLTonGPU *getGPUProductAsync(cuda::stream_t<>& cudaStream) const;
+
+private:
+  const SiPixelGainCalibrationForHLT *gains_ = nullptr;
+  SiPixelGainForHLTonGPU *gainForHLTonHost_ = nullptr;
+  struct GPUData {
+    ~GPUData();
+    SiPixelGainForHLTonGPU *gainForHLTonGPU = nullptr;
+    SiPixelGainForHLTonGPU_DecodingStructure *gainDataOnGPU = nullptr;
+  };
+  CUDAESProduct<GPUData> gpuData_;
+};
+
+#endif

CalibTracker/SiPixelESProducers/plugins/BuildFile.xml

@@ -6,6 +6,8 @@
 <use   name="Geometry/Records"/>
 <use   name="Geometry/TrackerGeometryBuilder"/>
 <use   name="CalibTracker/SiPixelESProducers"/>
+<use   name="HeterogeneousCore/CUDACore"/>
+<use   name="cuda-api-wrappers"/>
 <library   file="*.cc" name="CalibTrackerSiPixelESProducersPlugins">
   <flags   EDM_PLUGIN="1"/>
 </library>

CalibTracker/SiPixelESProducers/plugins/SiPixelGainCalibrationForHLTGPUESProducer.cc

@@ -0,0 +1,47 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+#include "CondFormats/DataRecord/interface/SiPixelGainCalibrationForHLTRcd.h"
+#include "FWCore/Framework/interface/ESProducer.h"
+#include "FWCore/Framework/interface/EventSetup.h"
+#include "FWCore/Framework/interface/ESHandle.h"
+#include "FWCore/Framework/interface/ModuleFactory.h"
+#include "FWCore/ParameterSet/interface/ParameterSet.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+
+#include <memory>
+
+class SiPixelGainCalibrationForHLTGPUESProducer: public edm::ESProducer {
+public:
+  explicit SiPixelGainCalibrationForHLTGPUESProducer(const edm::ParameterSet& iConfig);
+  std::unique_ptr<SiPixelGainCalibrationForHLTGPU> produce(const SiPixelGainCalibrationForHLTGPURcd& iRecord);
+
+  static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
+private:
+};
+
+SiPixelGainCalibrationForHLTGPUESProducer::SiPixelGainCalibrationForHLTGPUESProducer(const edm::ParameterSet& iConfig) {
+  setWhatProduced(this);
+}
+
+void SiPixelGainCalibrationForHLTGPUESProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
+  edm::ParameterSetDescription desc;
+  descriptions.add("siPixelGainCalibrationForHLTGPU", desc);
+}
+
+std::unique_ptr<SiPixelGainCalibrationForHLTGPU> SiPixelGainCalibrationForHLTGPUESProducer::produce(const SiPixelGainCalibrationForHLTGPURcd& iRecord) {
+  edm::ESHandle<SiPixelGainCalibrationForHLT> gains;
+  iRecord.getRecord<SiPixelGainCalibrationForHLTRcd>().get(gains);
+
+  edm::ESHandle<TrackerGeometry> geom;
+  iRecord.getRecord<TrackerDigiGeometryRecord>().get(geom);
+
+  return std::make_unique<SiPixelGainCalibrationForHLTGPU>(*gains, *geom);
+}
+
+#include "FWCore/Framework/interface/MakerMacros.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+
+DEFINE_FWK_EVENTSETUP_MODULE(SiPixelGainCalibrationForHLTGPUESProducer);

CalibTracker/SiPixelESProducers/src/ES_SiPixelGainCalibrationForHLTGPU.cc

@@ -0,0 +1,4 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+TYPELOOKUP_DATA_REG(SiPixelGainCalibrationForHLTGPU);

CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc

@@ -0,0 +1,98 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "Geometry/CommonDetUnit/interface/GeomDetType.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+
+#include <cuda.h>
+
+SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGainCalibrationForHLT& gains, const TrackerGeometry& geom):
+  gains_(&gains)
+{
+  // bizzarre logic (looking for fist strip-det) don't ask
+  auto const & dus = geom.detUnits();
+  unsigned m_detectors = dus.size();
+  for(unsigned int i=1;i<7;++i) {
+    if(geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]) != dus.size() &&
+        dus[geom.offsetDU(GeomDetEnumerators::tkDetEnum[i])]->type().isTrackerStrip()) {
+      if(geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]) < m_detectors) m_detectors = geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]);
+    }
+  }
+
+  /*
+  std::cout << "caching calibs for " << m_detectors << " pixel detectors of size " << gains.data().size() << std::endl;
+  std::cout << "sizes " << sizeof(char) << ' ' << sizeof(uint8_t) << ' ' << sizeof(SiPixelGainForHLTonGPU::DecodingStructure) << std::endl;
+  */
+
+  cudaCheck(cudaMallocHost((void**) & gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU)));
+  //gainForHLTonHost_->v_pedestals = gainDataOnGPU_; // how to do this?
+
+  // do not read back from the (possibly write-combined) memory buffer
+  auto minPed  = gains.getPedLow();
+  auto maxPed  = gains.getPedHigh();
+  auto minGain = gains.getGainLow();
+  auto maxGain = gains.getGainHigh();
+  auto nBinsToUseForEncoding = 253;
+
+  // we will simplify later (not everything is needed....)
+  gainForHLTonHost_->minPed_ = minPed;
+  gainForHLTonHost_->maxPed_ = maxPed;
+  gainForHLTonHost_->minGain_= minGain;
+  gainForHLTonHost_->maxGain_= maxGain;
+
+  gainForHLTonHost_->numberOfRowsAveragedOver_ = 80;
+  gainForHLTonHost_->nBinsToUseForEncoding_    = nBinsToUseForEncoding;
+  gainForHLTonHost_->deadFlag_                 = 255;
+  gainForHLTonHost_->noisyFlag_                = 254;
+
+  gainForHLTonHost_->pedPrecision  = static_cast<float>(maxPed - minPed) / nBinsToUseForEncoding;
+  gainForHLTonHost_->gainPrecision = static_cast<float>(maxGain - minGain) / nBinsToUseForEncoding;
+
+  /*
+  std::cout << "precisions g " << gainForHLTonHost_->pedPrecision << ' ' << gainForHLTonHost_->gainPrecision << std::endl;
+  */
+
+  // fill the index map
+  auto const & ind = gains.getIndexes();
+  /*
+  std::cout << ind.size() << " " << m_detectors << std::endl;
+  */
+
+  for (auto i=0U; i<m_detectors; ++i) {
+    auto p = std::lower_bound(ind.begin(),ind.end(),dus[i]->geographicalId().rawId(),SiPixelGainCalibrationForHLT::StrictWeakOrdering());
+    assert (p!=ind.end() && p->detid==dus[i]->geographicalId());
+    assert(p->iend<=gains.data().size());
+    assert(p->iend>=p->ibegin);
+    assert(0==p->ibegin%2);
+    assert(0==p->iend%2);
+    assert(p->ibegin!=p->iend);
+    assert(p->ncols>0);
+    gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(p->ibegin,p->iend), p->ncols);
+    // if (ind[i].detid!=dus[i]->geographicalId()) std::cout << ind[i].detid<<"!="<<dus[i]->geographicalId() << std::endl;
+    // gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(ind[i].ibegin,ind[i].iend), ind[i].ncols);
+  }
+
+}
+
+SiPixelGainCalibrationForHLTGPU::~SiPixelGainCalibrationForHLTGPU() {
+  cudaCheck(cudaFreeHost(gainForHLTonHost_));
+}
+
+SiPixelGainCalibrationForHLTGPU::GPUData::~GPUData() {
+  cudaCheck(cudaFree(gainForHLTonGPU));
+  cudaCheck(cudaFree(gainDataOnGPU));
+}
+
+const SiPixelGainForHLTonGPU *SiPixelGainCalibrationForHLTGPU::getGPUProductAsync(cuda::stream_t<>& cudaStream) const {
+  const auto& data = gpuData_.dataForCurrentDeviceAsync(cudaStream, [this](GPUData& data, cuda::stream_t<>& stream) {
+      cudaCheck(cudaMalloc((void**) & data.gainForHLTonGPU, sizeof(SiPixelGainForHLTonGPU)));
+      cudaCheck(cudaMalloc((void**) & data.gainDataOnGPU, this->gains_->data().size())); // TODO: this could be changed to cuda::memory::device::unique_ptr<>
+      // gains.data().data() is used also for non-GPU code, we cannot allocate it on aligned and write-combined memory
+      cudaCheck(cudaMemcpyAsync(data.gainDataOnGPU, this->gains_->data().data(), this->gains_->data().size(), cudaMemcpyDefault, stream.id()));
+
+      cudaCheck(cudaMemcpyAsync(data.gainForHLTonGPU, this->gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU), cudaMemcpyDefault, stream.id()));
+      cudaCheck(cudaMemcpyAsync(&(data.gainForHLTonGPU->v_pedestals), &(data.gainDataOnGPU), sizeof(SiPixelGainForHLTonGPU_DecodingStructure*), cudaMemcpyDefault, stream.id()));
+    });
+  return data.gainForHLTonGPU;
+}

CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h

@@ -0,0 +1,73 @@
+#ifndef CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
+#define CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
+
+#include <cstdint>
+#include <cstdio>
+#include <tuple>
+
+#include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
+
+struct SiPixelGainForHLTonGPU_DecodingStructure{
+  uint8_t gain;
+  uint8_t ped;
+};
+
+
+// copy of SiPixelGainCalibrationForHLT
+class SiPixelGainForHLTonGPU {
+
+ public:
+
+  using DecodingStructure = SiPixelGainForHLTonGPU_DecodingStructure;
+
+  using Range = std::pair<uint32_t,uint32_t>;
+
+
+  inline __host__ __device__
+  std::pair<float,float> getPedAndGain(uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn ) const {
+
+
+    auto range = rangeAndCols[moduleInd].first;
+    auto nCols = rangeAndCols[moduleInd].second;
+
+    // determine what averaged data block we are in (there should be 1 or 2 of these depending on if plaquette is 1 by X or 2 by X
+    unsigned int lengthOfColumnData  = (range.second-range.first)/nCols;
+    unsigned int lengthOfAveragedDataInEachColumn = 2;  // we always only have two values per column averaged block 
+    unsigned int numberOfDataBlocksToSkip = row / numberOfRowsAveragedOver_;
+
+
+    auto offset = range.first + col*lengthOfColumnData + lengthOfAveragedDataInEachColumn*numberOfDataBlocksToSkip;
+
+    assert(offset<range.second);
+    assert(offset<3088384);
+    assert(0==offset%2);
+
+    DecodingStructure const * __restrict__ lp = v_pedestals;
+    auto s = lp[offset/2];
+
+    isDeadColumn = (s.ped & 0xFF) == deadFlag_;
+    isNoisyColumn = (s.ped & 0xFF) == noisyFlag_;
+
+    return std::make_pair(decodePed(s.ped & 0xFF),decodeGain(s.gain & 0xFF));
+
+  }
+
+
+
+  constexpr float decodeGain(unsigned int gain) const {return gain*gainPrecision + minGain_;}
+  constexpr float decodePed (unsigned int ped) const { return ped*pedPrecision + minPed_;}
+
+  DecodingStructure * v_pedestals;
+  std::pair<Range, int> rangeAndCols[2000];
+
+  float  minPed_, maxPed_, minGain_, maxGain_;
+
+  float pedPrecision, gainPrecision;
+
+  unsigned int numberOfRowsAveragedOver_; // this is 80!!!!
+  unsigned int nBinsToUseForEncoding_;
+  unsigned int deadFlag_;
+  unsigned int noisyFlag_;
+};
+
+#endif // CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h

Configuration/Applications/python/ConfigBuilder.py

@@ -917,6 +917,8 @@ def define_Configs(self):
 		    self.loadAndRemember('SimGeneral.HepPDTESSource.'+self._options.particleTable+'_cfi')
 
         self.loadAndRemember('FWCore/MessageService/MessageLogger_cfi')
+        # Eventually replace with some more generic file to load
+        self.loadAndRemember('HeterogeneousCore/CUDAServices/CUDAService_cfi')
 
         self.ALCADefaultCFF="Configuration/StandardSequences/AlCaRecoStreams_cff"
         self.GENDefaultCFF="Configuration/StandardSequences/Generator_cff"

Configuration/ProcessModifiers/python/gpu_cff.py

@@ -0,0 +1,5 @@
+import FWCore.ParameterSet.Config as cms
+
+# This modifier is for replacing CPU modules with GPU counterparts
+
+gpu =  cms.Modifier()

Configuration/ProcessModifiers/python/riemannFitGPU_cff.py

@@ -0,0 +1,5 @@
+import FWCore.ParameterSet.Config as cms
+
+# This modifier is for replacing the default pixel track "fitting" with Riemann fit on GPU
+
+riemannFitGPU =  cms.Modifier()

Configuration/ProcessModifiers/python/riemannFit_cff.py

@@ -0,0 +1,5 @@
+import FWCore.ParameterSet.Config as cms
+
+# This modifier is for replacing the default pixel track "fitting" with Riemann fit
+
+riemannFit =  cms.Modifier()

Configuration/PyReleaseValidation/python/relval_2017.py

@@ -27,7 +27,7 @@
            10024.1,10024.2,10024.3,10024.4,10024.5,
            10801.0,10802.0,10803.0,10804.0,10805.0,10806.0,10807.0,10808.0,10809.0,10859.0,10871.0,
            10842.0,10824.0,10825.0,10826.0,10823.0,11024.0,11025.0,11224.0,
-           10824.1,10824.5,
+           10824.1,10824.5,10824.7,10824.8,10824.9,
            10824.6,11024.6,11224.6,
            11642.0,11624.0,11625.0,11626.0,11623.0,11824.0,11825.0,12024.0]
 for numWF in numWFIB:

Configuration/PyReleaseValidation/python/relval_standard.py

@@ -347,6 +347,9 @@
 workflows[136.862] = ['',['RunEGamma2018B','HLTDR2_2018','RECODR2_2018reHLT_skimEGamma_Prompt_L1TEgDQM','HARVEST2018_L1TEgDQM']]
 workflows[136.863] = ['',['RunDoubleMuon2018B','HLTDR2_2018','RECODR2_2018reHLT_Prompt','HARVEST2018']]
 workflows[136.864] = ['',['RunJetHT2018B','HLTDR2_2018','RECODR2_2018reHLT_skimJetHT_Prompt','HARVEST2018']]
+workflows[136.8645] = ['',['RunJetHT2018B','HLTDR2_2018','RECODR2_2018reHLT_Prompt_pixelTrackingOnly','HARVEST2018_pixelTrackingOnly']]
+workflows[136.8648] = ['',['RunJetHT2018B','HLTDR2_2018','RECODR2_2018reHLT_Prompt_pixelTrackingOnlyGPU','HARVEST2018_pixelTrackingOnly']]
+
 workflows[136.865] = ['',['RunMET2018B','HLTDR2_2018','RECODR2_2018reHLT_skimMET_Prompt','HARVEST2018']]
 workflows[136.866] = ['',['RunMuonEG2018B','HLTDR2_2018','RECODR2_2018reHLT_skimMuonEG_Prompt','HARVEST2018']]
 workflows[136.867] = ['',['RunSingleMu2018B','HLTDR2_2018','RECODR2_2018reHLT_skimSingleMu_Prompt_Lumi','HARVEST2018_L1TMuDQM']]

Configuration/PyReleaseValidation/python/relval_steps.py

@@ -1775,6 +1775,15 @@ def gen2018HiMix(fragment,howMuch):
     '--datatier': 'GEN-SIM-RECO,DQMIO',
     '--eventcontent': 'RECOSIM,DQM',
 }
+step3_riemannFit = {
+    '--procModifiers': 'riemannFit',
+}
+step3_riemannFitGPU = {
+    '--procModifiers': 'riemannFitGPU',
+}
+step3_gpu = {
+    '--procModifiers': 'gpu',
+}
 step3_trackingLowPU = {
     '--era': 'Run2_2016_trackingLowPU'
 }
@@ -1896,6 +1905,8 @@ def gen2018HiMix(fragment,howMuch):
 steps['RECODR2_2017reHLTSiPixelCalZeroBias_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2017reHLTSiPixelCalZeroBias']])
 
 steps['RECODR2_2018reHLT_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT']])
+steps['RECODR2_2018reHLT_Prompt_pixelTrackingOnly']=merge([{'-s': 'RAW2DIGI:RawToDigi_pixelOnly,RECO:reconstruction_pixelTrackingOnly,DQM:@pixelTrackingOnlyDQM'},steps['RECODR2_2018reHLT_Prompt']])
+steps['RECODR2_2018reHLT_Prompt_pixelTrackingOnlyGPU']=merge([step3_gpu, steps['RECODR2_2018reHLT_Prompt_pixelTrackingOnly']])
 steps['RECODR2_2018reHLT_skimEGamma_Prompt_L1TEgDQM']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT_skimEGamma_L1TEgDQM']])
 steps['RECODR2_2018reHLT_skimJetHT_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT_skimJetHT']])
 steps['RECODR2_2018reHLT_skimDisplacedJet_Prompt']=merge([{'--conditions':'auto:run2_data_promptlike'},steps['RECODR2_2018reHLT_skimDisplacedJet']])
@@ -2182,6 +2193,7 @@ def gen2018HiMix(fragment,howMuch):
 steps['HARVEST2018'] = merge([ {'--conditions':'auto:run2_data_relval','--era':'Run2_2018','--conditions':'auto:run2_data_promptlike',}, steps['HARVESTD'] ])
 steps['HARVEST2018_L1TEgDQM'] = merge([ {'-s':'HARVESTING:@standardDQM+@ExtraHLT+@miniAODDQM+@L1TEgamma'}, steps['HARVEST2018'] ])
 steps['HARVEST2018_L1TMuDQM'] = merge([ {'-s':'HARVESTING:@standardDQM+@ExtraHLT+@miniAODDQM+@L1TMuon'}, steps['HARVEST2018'] ])
+steps['HARVEST2018_pixelTrackingOnly'] = merge([ {'-s':'HARVESTING:@pixelTrackingOnlyDQM'}, steps['HARVEST2018'] ])
 steps['HARVEST2018_hBStar'] = merge([ {'--era' : 'Run2_2018_highBetaStar'}, steps['HARVEST2018'] ])
 steps['HARVEST2018_HEfail'] = merge([ {'--conditions':'auto:run2_data_promptlike_HEfail'}, steps['HARVEST2018'] ])
 steps['HARVEST2018_BadHcalMitig'] = merge([ {'--era' : 'Run2_2018,pf_badHcalMitigation','--conditions':'auto:run2_data_promptlike_HEfail'}, steps['HARVEST2018'] ])
@@ -2792,6 +2804,21 @@ def gen2018HiMix(fragment,howMuch):
         if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_pixelTrackingOnly, upgradeStepDict[step][k]])
         elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
 
+    for step in upgradeSteps['pixelTrackingOnlyRiemannFit']['steps']:
+        stepName = step + upgradeSteps['pixelTrackingOnlyRiemannFit']['suffix']
+        if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_riemannFit, step3_pixelTrackingOnly, upgradeStepDict[step][k]])
+        elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+
+    for step in upgradeSteps['pixelTrackingOnlyRiemannFitGPU']['steps']:
+        stepName = step + upgradeSteps['pixelTrackingOnlyRiemannFitGPU']['suffix']
+        if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_riemannFitGPU, step3_pixelTrackingOnly, upgradeStepDict[step][k]])
+        elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+
+    for step in upgradeSteps['pixelTrackingOnlyGPU']['steps']:
+        stepName = step + upgradeSteps['pixelTrackingOnlyGPU']['suffix']
+        if 'Reco' in step: upgradeStepDict[stepName][k] = merge([step3_gpu, step3_pixelTrackingOnly, upgradeStepDict[step][k]])
+        elif 'HARVEST' in step: upgradeStepDict[stepName][k] = merge([{'-s': 'HARVESTING:@trackingOnlyValidation+@pixelTrackingOnlyDQM'}, upgradeStepDict[step][k]])
+
     for step in upgradeSteps['trackingRun2']['steps']:
         stepName = step + upgradeSteps['trackingRun2']['suffix']
         if 'Reco' in step and upgradeStepDict[step][k]['--era']=='Run2_2017':