CUDA Vector for easier pushback

CondFormats/SiPixelObjects/src/SiPixelGainCalibrationOffline.cc

@@ -188,7 +188,7 @@ float SiPixelGainCalibrationOffline::getGain(const int& col, const int& row, con
   int maxRow = lengthOfColumnData - (lengthOfColumnData % numberOfRowsToAverageOver_) - 1;
   if (col >= nCols || row > maxRow){
     throw cms::Exception("CorruptedData")
-      << "[SiPixelGainCalibrationOffline::getPed] Pixel out of range: col " << col;
+      << "[SiPixelGainCalibrationOffline::getPed] Pixel out of range: col " << col << " row " << row;
   }  
   return decodeGain(s.datum & 0xFF);
 }

Configuration/DataProcessing/python/Merge.py

@@ -10,6 +10,7 @@
 
 from FWCore.ParameterSet.Config import Process, EndPath
 from FWCore.ParameterSet.Modules import OutputModule, Source, Service
+from Configuration.EventContent.EventContent_cff import NANOAODEventContent
 import FWCore.ParameterSet.Types as CfgTypes
 
 
@@ -64,7 +65,7 @@ def mergeProcess(*inputFiles, **options):
     if newDQMIO:
         outMod = OutputModule("DQMRootOutputModule")
     elif mergeNANO:
-        outMod = OutputModule("NanoAODOutputModule")
+        outMod = OutputModule("NanoAODOutputModule",NANOAODEventContent.clone())
     else:
         outMod = OutputModule("PoolOutputModule")
 

DQMOffline/Configuration/python/autoDQM.py

@@ -41,9 +41,9 @@
             'express':       ['@commonSiStripZeroBias+@muon+@hcal+@jetmet+@ecal',
                               'PostDQMOffline',
                               '@commonSiStripZeroBias+@muon+@hcal+@jetmet+@ecal'],
-            'allForPrompt':  ['@common+@muon+@hcal+@jetmet+@ecal',
+            'allForPrompt':  ['@common+@muon+@hcal+@jetmet+@ecal+@egamma',
                               'PostDQMOffline',
-                              '@common+@muon+@hcal+@jetmet+@ecal'],
+                              '@common+@muon+@hcal+@jetmet+@ecal+@egamma'],
             'miniAODDQM': ['DQMOfflineMiniAOD',
                            'PostDQMOfflineMiniAOD',
                            'DQMHarvestMiniAOD'],

DataFormats/PatCandidates/interface/Jet.h

@@ -47,6 +47,8 @@
 #include "DataFormats/Common/interface/OwnVector.h"
 #include "DataFormats/Common/interface/AtomicPtrCache.h"
 
+#include <numeric>
+
 
 // Define typedefs for convenience
 namespace pat {
@@ -498,21 +500,37 @@ namespace pat {
 
 
       /// String access to subjet list
-      pat::JetPtrCollection const & subjets( std::string label ) const ;
+      pat::JetPtrCollection const & subjets( std::string const & label ) const ;
 
       /// Add new set of subjets
-      void addSubjets( pat::JetPtrCollection const & pieces, std::string label = ""  );
+      void addSubjets( pat::JetPtrCollection const & pieces, std::string const & label = ""  );
 
       /// Check to see if the subjet collection exists
-      bool hasSubjets( std::string label ) const { return find( subjetLabels_.begin(), subjetLabels_.end(), label) != subjetLabels_.end(); }
+      bool hasSubjets( std::string const & label ) const { return find( subjetLabels_.begin(), subjetLabels_.end(), label) != subjetLabels_.end(); }
 
       /// Number of subjet collections
       unsigned int nSubjetCollections(  ) const { return  subjetCollections_.size(); }
 
       /// Subjet collection names
       std::vector<std::string> const & subjetCollectionNames() const { return subjetLabels_; }
 
-
+      /// Access to mass of subjets
+      double groomedMass(unsigned int index = 0) const{
+	auto const& sub = subjets(index);
+	return nSubjetCollections() > index && !sub.empty() ?
+	  std::accumulate( sub.begin(), sub.end(),
+			   reco::Candidate::LorentzVector(),
+			   [] (reco::Candidate::LorentzVector const & a, reco::CandidatePtr const & b){return a + b->p4();}).mass() :
+	  -1.0;
+      }
+      double groomedMass(std::string const & label) const{
+	auto const& sub = subjets(label);
+	return hasSubjets(label) && !sub.empty() ?
+	  std::accumulate( sub.begin(), sub.end(),
+			   reco::Candidate::LorentzVector(),
+			   [] (reco::Candidate::LorentzVector const & a, reco::CandidatePtr const & b){return a + b->p4();}).mass() :
+	  -1.0;
+      }
 
     protected:
 

DataFormats/PatCandidates/src/Jet.cc

@@ -599,7 +599,7 @@ pat::JetPtrCollection const & Jet::subjets( unsigned int index) const {
 
 
 /// String access to subjet list
-pat::JetPtrCollection const & Jet::subjets( std::string label ) const { 
+pat::JetPtrCollection const & Jet::subjets( std::string const & label ) const { 
   auto found = find( subjetLabels_.begin(), subjetLabels_.end(), label );
   if ( found != subjetLabels_.end() ){
     auto index = std::distance( subjetLabels_.begin(), found );
@@ -611,7 +611,7 @@ pat::JetPtrCollection const & Jet::subjets( std::string label ) const {
 }
 
 /// Add new set of subjets
-void Jet::addSubjets( pat::JetPtrCollection const & pieces, std::string label  ) {
+void Jet::addSubjets( pat::JetPtrCollection const & pieces, std::string const & label  ) {
   subjetCollections_.push_back( pieces );
   subjetLabels_.push_back( label );
 }

DataFormats/SiPixelDigi/interface/PixelDigi.h

@@ -17,7 +17,7 @@ class PixelDigi {
   typedef unsigned int PackedDigiType;
   typedef unsigned int ChannelType;
 
-  PixelDigi( int packed_value) : theData(packed_value) {}
+  explicit PixelDigi(PackedDigiType packed_value) : theData(packed_value) {}
 
   PixelDigi( int row, int col, int adc) {
     init( row, col, adc);

EventFilter/SiPixelRawToDigi/interface/GPUSimpleVector.hpp

@@ -0,0 +1,72 @@
+//  author: Felice Pantaleo, CERN, 2018
+#ifndef GPU_SIMPLEVECTOR_HPP_
+#define GPU_SIMPLEVECTOR_HPP_
+
+namespace GPU {
+template <class T> struct SimpleVector {
+  // Constructors
+  __host__ __device__ SimpleVector(unsigned int maxSize, T *m_data = nullptr)
+      : m_size(0), m_data(m_data), maxSize(static_cast<int>(maxSize)) {}
+
+  __host__ __device__ SimpleVector() : SimpleVector(0) {}
+
+  __inline__ __host__ __device__ int push_back(const T &element) {
+
+    auto previousSize = m_size;
+    m_size++;
+    if (previousSize < maxSize) {
+      m_data[previousSize] = element;
+      return previousSize;
+    } else {
+      --m_size;
+      return -1;
+    }
+  }
+
+#if defined(__NVCC__) || defined(__CUDACC__)
+  __device__ int push_back_ts(const T &element) {
+    auto previousSize = atomicAdd(&m_size, 1);
+    if (previousSize < maxSize) {
+      m_data[previousSize] = element;
+      return previousSize;
+    } else {
+      atomicSub(&m_size, 1);
+      return -1;
+    }
+  }
+#endif
+
+  __inline__ __host__ __device__ T pop_back() {
+    if (m_size > 0) {
+      auto previousSize = m_size--;
+      return m_data[previousSize - 1];
+    } else
+      return T();
+  }
+
+  __inline__ __host__ __device__ T operator[](int i) { return m_data[i]; }
+
+  __inline__ __host__ __device__ T at(int i) {
+    if (i < m_size)
+      return m_data[i];
+    else
+      return T();
+  }
+
+  __inline__ __host__ __device__ void reset() { m_size = 0; }
+
+  __inline__ __host__ __device__ int size() const { return m_size; }
+
+  __inline__ __host__ __device__ int capacity() const { return maxSize; }
+
+
+
+private:
+  int m_size;
+  int maxSize;
+
+  T *m_data;
+};
+} // namespace GPU
+
+#endif

EventFilter/SiPixelRawToDigi/plugins/BuildFile.xml

@@ -1,4 +1,9 @@
 <use   name="EventFilter/SiPixelRawToDigi"/>
-<library   file="*.cc" name="EventFilterSiPixelRawToDigiPlugins">
+<library   file="SiPixelDigiToRaw.cc SiPixelRawToDigi.cc SealModule.cc" name="EventFilterSiPixelRawToDigiPlugins">
   <flags   EDM_PLUGIN="1"/>
 </library>
+<library   file="SiPixelRawToDigiGPU.cc SiPixelFedCablingMapGPU.cc RawToDigiGPU.cu" name="EventFilterSiPixelRawToDigiGPUPlugins">
+  <use   name="cuda"/>
+  <flags   EDM_PLUGIN="1"/>
+  <flags   CUDA_FLAGS="-O2 --expt-relaxed-constexpr"/>
+</library>

EventFilter/SiPixelRawToDigi/plugins/DetParamBits.h

@@ -0,0 +1,39 @@
+// Sushil Dubey, Shashi Dugad, TIFR
+#ifndef DETPARAMBITS_H
+#define DETPARAMBITS_H
+typedef unsigned int uint;
+//reference
+//http://cmsdoxygen.web.cern.ch/cmsdoxygen/CMSSW_9_2_0/doc/html/d3/db2/PixelROC_8cc_source.html#l00197 
+const uint layerStartBit_  = 20;
+const uint ladderStartBit_ = 12;
+const uint moduleStartBit_ = 2;
+
+const uint panelStartBit_  = 10;
+const uint diskStartBit_   = 18;
+const uint bladeStartBit_  = 12;
+
+const uint layerMask_      = 0xF;
+const uint ladderMask_     = 0xFF;
+const uint moduleMask_     = 0x3FF;
+const uint panelMask_      = 0x3;
+const uint diskMask_       = 0xF;
+const uint bladeMask_      = 0x3F;
+
+// __host__ __device__ bool isBarrel(uint rawId) {
+//   return (1==((rawId>>25)&0x7));
+// }
+
+__host__ __device__ int getLayer(uint rawId) {
+  int layer  = (rawId >> layerStartBit_)  & layerMask_;
+  return layer;
+}
+
+__host__ __device__ int getDisk(uint rawId) {
+  // int side =1;
+  // unsigned int panel = ((rawId>>panelStartBit_) & panelMask_);
+  // if(panel==1) side = -1;
+   unsigned int disk = int((rawId>>diskStartBit_) & diskMask_);
+  // return disk*side;
+  return disk;
+}
+#endif

EventFilter/SiPixelRawToDigi/plugins/EventInfoGPU.h

@@ -0,0 +1,11 @@
+/*Sushil Dubey, Shashi Dugad, TIFR
+*/
+
+#ifndef EVENTINFO_GPU
+#define EVENTINFO_GPU
+
+const int NEVENT = 1 ;  //optimal number of events to run simultaneously,
+// using 4 cuda stream, hence it should be multiple of 4
+const int NMODULE = 1856; // for phase 1, we have 1856 modules
+
+#endif