Phase II Patatrack Pixel Local Reco

CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h

@@ -11,11 +11,13 @@ namespace gpuClustering {
 #else
   // optimized for real data PU 50
   // tested on MC events with 55-75 pileup events
-  constexpr uint32_t maxHitsInIter() { return 160; }
+  constexpr uint32_t maxHitsInIter() { return 160; }  //TODO better tuning for PU 140-200
 #endif
   constexpr uint32_t maxHitsInModule() { return 1024; }
 
-  constexpr uint16_t maxNumModules = 2000;
+  constexpr uint32_t maxNumDigis = 3 * 256 * 1024;  // @PU=200 µ=530k σ=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;

CUDADataFormats/Track/interface/TrackSoAHeterogeneousT.h

@@ -5,7 +5,7 @@
 #include <algorithm>
 
 #include "CUDADataFormats/Track/interface/TrajectoryStateSoAT.h"
-#include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h"
+#include "Geometry/TrackerGeometryBuilder/interface/pixelTopology.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h"
 
 #include "CUDADataFormats/Common/interface/HeterogeneousSoA.h"

CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h

@@ -3,6 +3,7 @@
 
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h"
 #include "CUDADataFormats/Common/interface/HeterogeneousSoA.h"
+#include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
 
 template <typename Traits>
 class TrackingRecHit2DHeterogeneous {
@@ -16,6 +17,7 @@ class TrackingRecHit2DHeterogeneous {
 
   explicit TrackingRecHit2DHeterogeneous(
       uint32_t nHits,
+      bool isPhase2,
       int32_t offsetBPIX2,
       pixelCPEforGPU::ParamsOnGPU const* cpeParams,
       uint32_t const* hitsModuleStart,
@@ -33,6 +35,7 @@ class TrackingRecHit2DHeterogeneous {
   TrackingRecHit2DSOAView const* view() const { return m_view.get(); }
 
   auto nHits() const { return m_nHits; }
+  auto nMaxModules() const { return m_nMaxModules; }
   auto offsetBPIX2() const { return m_offsetBPIX2; }
 
   auto hitsModuleStart() const { return m_hitsModuleStart; }
@@ -66,6 +69,7 @@ class TrackingRecHit2DHeterogeneous {
 
   uint32_t const* m_hitsModuleStart;  // needed for legacy, this is on GPU!
 
+  uint32_t m_nMaxModules;
   // needed as kernel params...
   PhiBinner* m_phiBinner;
   PhiBinner::index_type* m_phiBinnerStorage;
@@ -83,6 +87,7 @@ using TrackingRecHit2DHost = TrackingRecHit2DHeterogeneous<cms::cudacompat::Host
 template <typename Traits>
 TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(
     uint32_t nHits,
+    bool isPhase2,
     int32_t offsetBPIX2,
     pixelCPEforGPU::ParamsOnGPU const* cpeParams,
     uint32_t const* hitsModuleStart,
@@ -91,7 +96,10 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(
     : m_nHits(nHits), m_offsetBPIX2(offsetBPIX2), m_hitsModuleStart(hitsModuleStart) {
   auto view = Traits::template make_host_unique<TrackingRecHit2DSOAView>(stream);
 
+  m_nMaxModules = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules;
+
   view->m_nHits = nHits;
+  view->m_nMaxModules = m_nMaxModules;
   m_view = Traits::template make_unique<TrackingRecHit2DSOAView>(stream);  // leave it on host and pass it by value?
   m_AverageGeometryStore = Traits::template make_unique<TrackingRecHit2DSOAView::AverageGeometry>(stream);
   view->m_averageGeometry = m_AverageGeometryStore.get();
@@ -120,8 +128,11 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(
     copyFromGPU(input, stream);
   } else {
     assert(input == nullptr);
+
+    auto nL = isPhase2 ? phase2PixelTopology::numberOfLayers : phase1PixelTopology::numberOfLayers;
+
     m_store16 = Traits::template make_unique<uint16_t[]>(nHits * n16, stream);
-    m_store32 = Traits::template make_unique<float[]>(nHits * n32 + phase1PixelTopology::numberOfLayers + 1, stream);
+    m_store32 = Traits::template make_unique<float[]>(nHits * n32 + nL + 1, stream);
     m_PhiBinnerStore = Traits::template make_unique<TrackingRecHit2DSOAView::PhiBinner>(stream);
   }
 

CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h

@@ -6,7 +6,7 @@
 #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
-#include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h"
+#include "Geometry/TrackerGeometryBuilder/interface/pixelTopology.h"
 #include "CUDADataFormats/TrackingRecHit/interface/SiPixelHitStatus.h"
 
 namespace pixelCPEforGPU {
@@ -20,15 +20,17 @@ class TrackingRecHit2DSOAView {
 
   using hindex_type = uint32_t;  // if above is <=2^32
 
-  using PhiBinner = cms::cuda::HistoContainer<int16_t, 128, -1, 8 * sizeof(int16_t), hindex_type, 10>;
+  using PhiBinner = cms::cuda::
+      HistoContainer<int16_t, 256, -1, 8 * sizeof(int16_t), hindex_type, pixelTopology::maxLayers>;  //28 for phase2 geometry
 
-  using AverageGeometry = phase1PixelTopology::AverageGeometry;
+  using AverageGeometry = pixelTopology::AverageGeometry;
 
   template <typename>
   friend class TrackingRecHit2DHeterogeneous;
   friend class TrackingRecHit2DReduced;
 
   __device__ __forceinline__ uint32_t nHits() const { return m_nHits; }
+  __device__ __forceinline__ uint32_t nMaxModules() const { return m_nMaxModules; }
 
   __device__ __forceinline__ float& xLocal(int i) { return m_xl[i]; }
   __device__ __forceinline__ float xLocal(int i) const { return __ldg(m_xl + i); }
@@ -114,6 +116,7 @@ class TrackingRecHit2DSOAView {
   PhiBinner::index_type* m_phiBinnerStorage;
 
   uint32_t m_nHits;
+  uint32_t m_nMaxModules;
 };
 
 #endif  // CUDADataFormats_TrackingRecHit_interface_TrackingRecHit2DSOAView_h

CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc

@@ -13,9 +13,9 @@ cms::cuda::host::unique_ptr<float[]> TrackingRecHit2DGPU::localCoordToHostAsync(
 
 template <>
 cms::cuda::host::unique_ptr<uint32_t[]> TrackingRecHit2DGPU::hitsModuleStartToHostAsync(cudaStream_t stream) const {
-  auto ret = cms::cuda::make_host_unique<uint32_t[]>(gpuClustering::maxNumModules + 1, stream);
-  cudaCheck(cudaMemcpyAsync(
-      ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClustering::maxNumModules + 1), cudaMemcpyDefault, stream));
+  auto ret = cms::cuda::make_host_unique<uint32_t[]>(nMaxModules() + 1, stream);
+  cudaCheck(
+      cudaMemcpyAsync(ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (nMaxModules() + 1), cudaMemcpyDefault, stream));
   return ret;
 }
 

CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp

@@ -16,16 +16,18 @@ int main() {
   cudaCheck(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
 
   auto nHits = 200;
+  auto nModules = 2000;
   // inner scope to deallocate memory before destroying the stream
   {
-    TrackingRecHit2DGPU tkhit(nHits, 0, nullptr, nullptr, stream);
+    TrackingRecHit2DGPU tkhit(nHits, nModules, 0, nullptr, nullptr, stream);
 
     testTrackingRecHit2D::runKernels(tkhit.view());
 
-    TrackingRecHit2DHost tkhitH(nHits, 0, nullptr, nullptr, stream, &tkhit);
+    TrackingRecHit2DHost tkhitH(nHits, nModules, 0, nullptr, nullptr, stream, &tkhit);
     cudaStreamSynchronize(stream);
     assert(tkhitH.view());
     assert(tkhitH.view()->nHits() == unsigned(nHits));
+    assert(tkhitH.view()->nMaxModules() == unsigned(nModules));
   }
 
   cudaCheck(cudaStreamDestroy(stream));

CalibTracker/SiPixelLorentzAngle/src/SiPixelLorentzAnglePCLWorker.cc

@@ -48,7 +48,7 @@
 #include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
 #include "Geometry/TrackerGeometryBuilder/interface/PixelTopologyMap.h"
 #include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
-#include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h"
+#include "Geometry/TrackerGeometryBuilder/interface/pixelTopology.h"
 #include "RecoTracker/TransientTrackingRecHit/interface/TkTransientTrackingRecHitBuilder.h"
 #include "TrackingTools/PatternTools/interface/TrajTrackAssociation.h"
 #include "TrackingTools/Records/interface/TransientRecHitRecord.h"

CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h

@@ -16,7 +16,7 @@
 #endif  // __device__
 #endif  // __CUDACC__
 
-#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
+#include "Geometry/TrackerGeometryBuilder/interface/pixelTopology.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
 
 struct SiPixelGainForHLTonGPU_DecodingStructure {
@@ -35,7 +35,6 @@ class SiPixelGainForHLTonGPU {
       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
@@ -60,7 +59,7 @@ class SiPixelGainForHLTonGPU {
   constexpr float decodePed(unsigned int ped) const { return float(ped) * pedPrecision_ + minPed_; }
 
   DecodingStructure* v_pedestals_;
-  std::pair<Range, int> rangeAndCols_[gpuClustering::maxNumModules];
+  std::pair<Range, int> rangeAndCols_[phase1PixelTopology::numberOfModules];
 
   float minPed_, maxPed_, minGain_, maxGain_;
   float pedPrecision_, gainPrecision_;

Configuration/PyReleaseValidation/python/relval_2026.py

@@ -4,7 +4,7 @@
 # here only define the workflows as a combination of the steps defined above:
 workflows = Matrix()
 
-# each workflow defines a name and a list of steps to be done. 
+# each workflow defines a name and a list of steps to be done.
 # if no explicit name/label given for the workflow (first arg),
 # the name of step1 will be used
 
@@ -39,7 +39,7 @@
 numWFIB.extend([38234.0]) #2026D85
 numWFIB.extend([38634.0]) #2026D86
 numWFIB.extend([39034.0]) #2026D87
-numWFIB.extend([39434.0]) #2026D88
+numWFIB.extend([39434.0,39434.5,39434.501,39434.502]) #2026D88, pixelTrackingOnly, Patatrack local reconstruction on CPU, Patatrack local reconstruction on GPU
 numWFIB.extend([39834.0]) #2026D89
 numWFIB.extend([40234.0]) #2026D90
 

Configuration/PyReleaseValidation/python/upgradeWorkflowComponents.py

@@ -156,7 +156,7 @@ def setup_(self, step, stepName, stepDict, k, properties):
         era=properties.get('Era', None)
         modifier=properties.get('ProcessModifier',None)
         if cust is not None: stepDict[stepName][k]['--customise']=cust
-        if era is not None: 
+        if era is not None:
             stepDict[stepName][k]['--era']=era
         if modifier is not None: stepDict[stepName][k]['--procModifier']=modifier
     def condition(self, fragment, stepList, key, hasHarvest):
@@ -304,7 +304,7 @@ def setup__(self, step, stepName, stepDict, k, properties):
         if 'Reco' in step: stepDict[stepName][k] = merge([self.step3, stepDict[step][k]])
         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
+        return '2017' in key or '2018' in key or '2021' in key or '2026' in key
 upgradeWFs['pixelTrackingOnly'] = UpgradeWorkflow_pixelTrackingOnly(
     steps = [
         'Reco',
@@ -482,6 +482,7 @@ def condition(self, fragment, stepList, key, hasHarvest):
             ('2021' in key and fragment == "TTbar_14TeV"),
             ('2018' in key and fragment == "ZMM_13"),
             ('2021' in key and fragment == "ZMM_14"),
+            ('2026D88' in key and fragment == "TTbar_14TeV" and "PixelOnly" in self.suffix)
         ]
         result = any(selected) and hasHarvest
 

DQMOffline/Configuration/python/DQMOffline_cff.py

@@ -168,11 +168,18 @@
                                         pixelPVMonitor *
                                         monitorpixelSoASource )
 
+
+from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker
+_DQMOfflinePixelTrackingNoSoA = DQMOfflinePixelTracking.copy()
+_DQMOfflinePixelTrackingNoSoA = cms.Sequence(pixelTracksMonitoring * pixelPVMonitor)
+
+phase2_tracker.toReplaceWith(DQMOfflinePixelTracking, _DQMOfflinePixelTrackingNoSoA)
+
 DQMOuterTracker = cms.Sequence( DQMOfflineDCS *
                                 OuterTrackerSource *
                                 DQMMessageLogger *
                                 DQMOfflinePhysics *
-                                DQMOfflineVertex 
+                                DQMOfflineVertex
                                 )
 
 DQMOfflineTrackerPhase2 = cms.Sequence( trackerphase2DQMSource )
@@ -186,7 +193,7 @@
 
 DQMOfflineCommon = cms.Sequence( DQMOfflineDCS *
                                  DQMMessageLogger *
-				 DQMOfflineTrackerStrip * 
+				 DQMOfflineTrackerStrip *
 				 DQMOfflineTrackerPixel *
                                  DQMOfflineTracking *
                                  DQMOfflineTrigger *

EventFilter/SiPixelRawToDigi/python/SiPixelRawToDigi_cfi.py

@@ -15,14 +15,14 @@
 )
 
 from Configuration.ProcessModifiers.siPixelQualityRawToDigi_cff import siPixelQualityRawToDigi
-siPixelQualityRawToDigi.toModify(siPixelDigis.cpu,  
-    UseQualityInfo = True,                         
+siPixelQualityRawToDigi.toModify(siPixelDigis.cpu,
+    UseQualityInfo = True,
     SiPixelQualityLabel = 'forRawToDigi',
 )
 
-
+from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker
 # SwitchProducer wrapping the legacy pixel digis producer or an alias combining the pixel digis information converted from SoA
-gpu.toModify(siPixelDigis,
+(gpu & ~phase2_tracker).toModify(siPixelDigis,
     cuda = cms.EDAlias(
         siPixelDigiErrors = cms.VPSet(
             cms.PSet(type = cms.string("DetIdedmEDCollection")),

EventFilter/SiPixelRawToDigi/python/siPixelDigis_cff.py

@@ -31,7 +31,8 @@
 
 
 from Configuration.ProcessModifiers.gpu_cff import gpu
-gpu.toReplaceWith(siPixelDigisTask, cms.Task(
+from Configuration.Eras.Modifier_phase2_tracker_cff import phase2_tracker
+(gpu & ~phase2_tracker).toReplaceWith(siPixelDigisTask, cms.Task(
     # copy the pixel digis (except errors) and clusters to the host
     siPixelDigisSoA,
     # copy the pixel digis errors to the host