Synchronise with CMSSW_10_2_0_pre5

cms-patatrack · Jun 7, 2018 · 655e4ed · 655e4ed
2 parents 30c7b03 + 571e8b8
commit 655e4ed
Show file tree

Hide file tree

Showing 113 changed files with 65,441 additions and 40 deletions.
diff --git a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h
@@ -0,0 +1,71 @@
+#ifndef CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
+#define CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h
+
+#include <cassert>
+#include <cstdint>
+#include <cstdio>
+#include <tuple>
+
+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);
+
+    auto s = v_pedestals[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
diff --git a/Configuration/Applications/python/ConfigBuilder.py b/Configuration/Applications/python/ConfigBuilder.py
@@ -916,6 +916,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"

diff --git a/Configuration/ProcessModifiers/python/gpu_cff.py b/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()
diff --git a/Configuration/ProcessModifiers/python/riemannFit_cff.py b/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()
diff --git a/Configuration/PyReleaseValidation/python/relval_2017.py b/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.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:

diff --git a/Configuration/PyReleaseValidation/python/relval_steps.py b/Configuration/PyReleaseValidation/python/relval_steps.py
@@ -1675,6 +1675,12 @@ def gen2018HiMix(fragment,howMuch):
     '--datatier': 'GEN-SIM-RECO,DQMIO',
     '--eventcontent': 'RECOSIM,DQM',
 }
+step3_riemannFit = {
+    '--procModifiers': 'riemannFit',
+}
+step3_gpu = {
+    '--procModifiers': 'gpu',
+}
 step3_trackingLowPU = {
     '--era': 'Run2_2016_trackingLowPU'
 }
@@ -2652,6 +2658,16 @@ 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['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':

diff --git a/Configuration/PyReleaseValidation/python/relval_upgrade.py b/Configuration/PyReleaseValidation/python/relval_upgrade.py
@@ -74,15 +74,18 @@ def makeStepName(key,frag,step,suffix):
             # special workflows for tracker
             if (upgradeDatasetFromFragment[frag]=="TTbar_13" or upgradeDatasetFromFragment[frag]=="TTbar_14TeV") and not 'PU' in key and hasHarvest:
                 # skip ALCA
-                trackingVariations = ['trackingOnly','trackingRun2','trackingOnlyRun2','trackingLowPU','pixelTrackingOnly']
+                trackingVariations = ['trackingOnly','trackingRun2','trackingOnlyRun2','trackingLowPU','pixelTrackingOnly','pixelTrackingOnlyRiemannFit','pixelTrackingOnlyGPU']
                 for tv in trackingVariations:
                     stepList[tv] = filter(lambda s : "ALCA" not in s, stepList[tv])
                 workflows[numWF+upgradeSteps['trackingOnly']['offset']] = [ upgradeDatasetFromFragment[frag], stepList['trackingOnly']]
                 if '2017' in key:
                     for tv in trackingVariations[1:]:
                         workflows[numWF+upgradeSteps[tv]['offset']] = [ upgradeDatasetFromFragment[frag], stepList[tv]]
                 elif '2018' in key:
-                    workflows[numWF+upgradeSteps['pixelTrackingOnly']['offset']] = [ upgradeDatasetFromFragment[frag], stepList['pixelTrackingOnly']]
+                    for tv in trackingVariations:
+                        if not "pixelTrackingOnly" in tv:
+                            continue
+                        workflows[numWF+upgradeSteps[tv]['offset']] = [ upgradeDatasetFromFragment[frag], stepList[tv]]
 
             # special workflows for HE
             if upgradeDatasetFromFragment[frag]=="TTbar_13" and '2018' in key:

diff --git a/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py b/Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py
@@ -140,6 +140,28 @@
     'suffix' : '_pixelTrackingOnly',
     'offset' : 0.5,
 }
+upgradeSteps['pixelTrackingOnlyRiemannFit'] = {
+    'steps' : [
+        'RecoFull',
+        'HARVESTFull',
+        'RecoFullGlobal',
+        'HARVESTFullGlobal',
+    ],
+    'PU' : [],
+    'suffix' : '_pixelTrackingOnlyRiemannFit',
+    'offset' : 0.7,
+}
+upgradeSteps['pixelTrackingOnlyGPU'] = {
+    'steps' : [
+        'RecoFull',
+        'HARVESTFull',
+        'RecoFullGlobal',
+        'HARVESTFullGlobal',
+    ],
+    'PU' : [],
+    'suffix' : '_pixelTrackingOnlyGPU',
+    'offset' : 0.8,
+}
 upgradeSteps['Timing'] = {
     'steps' : upgradeSteps['baseline']['steps'],
     'PU' : upgradeSteps['baseline']['PU'],

diff --git a/Configuration/StandardSequences/python/RawToDigi_cff.py b/Configuration/StandardSequences/python/RawToDigi_cff.py
@@ -4,6 +4,7 @@
 # scenarios. In this case it makes changes for Run 2.
 
 from EventFilter.SiPixelRawToDigi.SiPixelRawToDigi_cfi import *
+from EventFilter.SiPixelRawToDigi.siPixelDigisHeterogeneous_cfi import *
 
 from EventFilter.SiStripRawToDigi.SiStripDigis_cfi import *
 
@@ -59,6 +60,10 @@
                          +tcdsDigis
                          +onlineMetaDataDigis
                          )
+from Configuration.ProcessModifiers.gpu_cff import gpu
+_RawToDigi_gpu = RawToDigi.copy()
+_RawToDigi_gpu.replace(siPixelDigis, siPixelDigisHeterogeneous+siPixelDigis)
+gpu.toReplaceWith(RawToDigi, _RawToDigi_gpu)
 
 RawToDigi_noTk = cms.Sequence(L1TRawToDigi
                               +ecalDigis
@@ -73,10 +78,11 @@
                               +onlineMetaDataDigis
                               )
 
-RawToDigi_pixelOnly = cms.Sequence(siPixelDigis)
+RawToDigi_pixelOnly = cms.Sequence(siPixelDigisHeterogeneous+siPixelDigis)
 
 scalersRawToDigi.scalersInputTag = 'rawDataCollector'
-siPixelDigis.InputLabel = 'rawDataCollector'
+siPixelDigisHeterogeneous.InputLabel = 'rawDataCollector'
+(~gpu).toModify(siPixelDigis, InputLabel = 'rawDataCollector')
 #false by default anyways ecalDigis.DoRegional = False
 ecalDigis.InputLabel = 'rawDataCollector'
 ecalPreshowerDigis.sourceTag = 'rawDataCollector'

diff --git a/Configuration/StandardSequences/python/Reconstruction_cff.py b/Configuration/StandardSequences/python/Reconstruction_cff.py
@@ -15,7 +15,7 @@
 
 siPixelClusterShapeCachePreSplitting = siPixelClusterShapeCache.clone(
     src = 'siPixelClustersPreSplitting'
-    )
+)
 
 # Global  reco
 from RecoEcal.Configuration.RecoEcal_cff import *

diff --git a/DataFormats/GeometrySurface/test/BuildFile.xml b/DataFormats/GeometrySurface/test/BuildFile.xml
@@ -14,3 +14,16 @@
 </bin>
 <bin file="Bounds_t.cpp">
 </bin>
+
+<bin file="gpuFrameTransformKernel.cu, gpuFrameTransformTest.cpp" name="gpuFrameTransformTest">
+  <use name="cuda"/>
+  <use name="cuda-api-wrappers"/>
+  <flags CXXFLAGS="-g"/>
+</bin>
+
+<bin file="gpuFrameTransformKernel.cu, gpuFrameTransformTest.cpp" name="gpuFrameTransformTestRep">
+  <use name="cuda"/>
+  <use name="cuda-api-wrappers"/>
+  <flags CXXFLAGS="-ffp-contract=off"/>
+  <flags CUDA_FLAGS="-fmad=false -ftz=false -prec-div=true -prec-sqrt=true"/>
+</bin>
diff --git a/DataFormats/GeometrySurface/test/gpuFrameTransformKernel.cu b/DataFormats/GeometrySurface/test/gpuFrameTransformKernel.cu
@@ -0,0 +1,43 @@
+#include "DataFormats/GeometrySurface/interface/SOARotation.h"
+#include <cstdint>
+
+
+__global__
+void toGlobal(SOAFrame<float> const * frame, 
+	      float const * xl, float const * yl,
+	      float * x, float * y, float * z,
+	      uint32_t n)
+{
+  int i = blockDim.x * blockIdx.x + threadIdx.x;
+  if (i >= n) return;
+
+  frame[0].toGlobal(xl[i],yl[i],x[i],y[i],z[i]);
+
+
+}
+
+#include<iostream>
+#include <iomanip>
+#include "cuda/api_wrappers.h"
+
+
+void toGlobalWrapper(SOAFrame<float> const * frame, 
+	      float const * xl, float const * yl,
+	      float * x, float * y, float * z,
+		     uint32_t n) {
+
+  int threadsPerBlock = 256;
+  int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
+  std::cout
+    << "CUDA toGlobal kernel launch with " << blocksPerGrid
+    << " blocks of " << threadsPerBlock << " threads" << std::endl;
+
+  cuda::launch(
+	       toGlobal,
+	       { blocksPerGrid, threadsPerBlock },
+	       frame, xl, yl, x, y, z,
+	       n
+	       );
+
+}
+