Clean up the pixel local reconstruction code

CUDADataFormats/Common/BuildFile.xml

@@ -1,7 +1,7 @@
 <iftool name="cuda">
+  <use name="rootcore"/>
   <use name="DataFormats/Common"/>
   <use name="HeterogeneousCore/CUDAUtilities"/>
-  <use name="rootcore"/>
   <export>
     <lib name="1"/>
   </export>

CUDADataFormats/SiPixelCluster/BuildFile.xml

@@ -1,10 +1,9 @@
-<use name="DataFormats/Common"/>
-<use name="CUDADataFormats/Common"/>
-<use name="HeterogeneousCore/CUDAUtilities"/>
 <use name="cuda"/>
 <use name="rootcore"/>
+<use name="CUDADataFormats/Common"/>
+<use name="DataFormats/Common"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
 
 <export>
     <lib name="1"/>
 </export>
-

CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h

@@ -60,4 +60,4 @@ class SiPixelClustersCUDA {
   uint32_t nClusters_h = 0;
 };
 
-#endif
+#endif  // CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h

CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h

@@ -2,30 +2,35 @@
 #define CUDADataFormats_SiPixelCluster_interface_gpuClusteringConstants_h
 
 #include <cstdint>
+#include <limits>
 
 namespace pixelGPUConstants {
 #ifdef GPU_SMALL_EVENTS
+  // kept for testing and debugging
   constexpr uint32_t maxNumberOfHits = 24 * 1024;
 #else
-  constexpr uint32_t maxNumberOfHits =
-      48 * 1024;  // data at pileup 50 has 18300 +/- 3500 hits; 40000 is around 6 sigma away
+  // data at pileup 50 has 18300 +/- 3500 hits; 40000 is around 6 sigma away
+  // tested on MC events with 55-75 pileup events
+  constexpr uint32_t maxNumberOfHits = 48 * 1024;
 #endif
 }  // namespace pixelGPUConstants
 
 namespace gpuClustering {
 #ifdef GPU_SMALL_EVENTS
+  // kept for testing and debugging
   constexpr uint32_t maxHitsInIter() { return 64; }
 #else
   // optimized for real data PU 50
+  // tested on MC events with 55-75 pileup events
   constexpr uint32_t maxHitsInIter() { return 160; }
 #endif
   constexpr uint32_t maxHitsInModule() { return 1024; }
 
-  constexpr uint32_t MaxNumModules = 2000;
-  constexpr int32_t MaxNumClustersPerModules = maxHitsInModule();
-  constexpr uint32_t MaxHitsInModule = maxHitsInModule();  // as above
-  constexpr uint32_t MaxNumClusters = pixelGPUConstants::maxNumberOfHits;
-  constexpr uint16_t InvId = 9999;  // must be > MaxNumModules
+  constexpr uint16_t maxNumModules = 2000;
+  constexpr int32_t maxNumClustersPerModules = maxHitsInModule();
+  constexpr uint32_t maxNumClusters = pixelGPUConstants::maxNumberOfHits;
+  constexpr uint16_t invalidModuleId = std::numeric_limits<uint16_t>::max() - 1;
+  static_assert(invalidModuleId > maxNumModules);  // invalidModuleId must be > maxNumModules
 
 }  // namespace gpuClustering
 

CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc

@@ -1,8 +1,7 @@
 #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
-
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 
 SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
     : moduleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)),

CUDADataFormats/SiPixelCluster/src/classes.h

@@ -1,8 +1,8 @@
-#ifndef CUDADataFormats_SiPixelCluster_classes_h
-#define CUDADataFormats_SiPixelCluster_classes_h
+#ifndef CUDADataFormats_SiPixelCluster_src_classes_h
+#define CUDADataFormats_SiPixelCluster_src_classes_h
 
 #include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
 #include "DataFormats/Common/interface/Wrapper.h"
 
-#endif
+#endif  // CUDADataFormats_SiPixelCluster_src_classes_h

CUDADataFormats/SiPixelDigi/BuildFile.xml

@@ -1,9 +1,9 @@
+<use name="cuda"/>
+<use name="rootcore"/>
 <use name="CUDADataFormats/Common"/>
 <use name="DataFormats/Common"/>
 <use name="DataFormats/SiPixelRawData"/>
 <use name="HeterogeneousCore/CUDAUtilities"/>
-<use name="cuda"/>
-<use name="rootcore"/>
 
 <export>
     <lib name="1"/>

CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h

@@ -1,13 +1,13 @@
 #ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
 #define CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
 
+#include <cuda_runtime.h>
+
 #include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h"
 #include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h"
-
-#include <cuda_runtime.h>
 
 class SiPixelDigiErrorsCUDA {
 public:
@@ -39,4 +39,4 @@ class SiPixelDigiErrorsCUDA {
   SiPixelFormatterErrors formatterErrors_h;
 };
 
-#endif
+#endif  // CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h

CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h

@@ -1,12 +1,12 @@
 #ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
 #define CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
 
+#include <cuda_runtime.h>
+
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
 
-#include <cuda_runtime.h>
-
 class SiPixelDigisCUDA {
 public:
   SiPixelDigisCUDA() = default;
@@ -82,4 +82,4 @@ class SiPixelDigisCUDA {
   uint32_t nDigis_h = 0;
 };
 
-#endif
+#endif  // CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h

CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc

@@ -1,12 +1,11 @@
-#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+#include <cassert>
 
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/memsetAsync.h"
 
-#include <cassert>
-
 SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream)
     : data_d(cms::cuda::make_device_unique<SiPixelErrorCompact[]>(maxFedWords, stream)),
       error_d(cms::cuda::make_device_unique<SiPixelErrorCompactVector>(stream)),

CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc

@@ -1,27 +1,23 @@
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
-
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
-
-SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream) {
-  xx_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  yy_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  adc_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  moduleInd_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  clus_d = cms::cuda::make_device_unique<int32_t[]>(maxFedWords, stream);
-
-  pdigi_d = cms::cuda::make_device_unique<uint32_t[]>(maxFedWords, stream);
-  rawIdArr_d = cms::cuda::make_device_unique<uint32_t[]>(maxFedWords, stream);
 
+SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream)
+    : xx_d(cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream)),
+      yy_d(cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream)),
+      adc_d(cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream)),
+      moduleInd_d(cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream)),
+      clus_d(cms::cuda::make_device_unique<int32_t[]>(maxFedWords, stream)),
+      view_d(cms::cuda::make_device_unique<DeviceConstView>(stream)),
+      pdigi_d(cms::cuda::make_device_unique<uint32_t[]>(maxFedWords, stream)),
+      rawIdArr_d(cms::cuda::make_device_unique<uint32_t[]>(maxFedWords, stream)) {
   auto view = cms::cuda::make_host_unique<DeviceConstView>(stream);
   view->xx_ = xx_d.get();
   view->yy_ = yy_d.get();
   view->adc_ = adc_d.get();
   view->moduleInd_ = moduleInd_d.get();
   view->clus_ = clus_d.get();
-
-  view_d = cms::cuda::make_device_unique<DeviceConstView>(stream);
   cms::cuda::copyAsync(view_d, view, stream);
 }
 

CUDADataFormats/SiPixelDigi/src/classes.h

@@ -1,9 +1,9 @@
-#ifndef CUDADataFormats_SiPixelDigi_classes_h
-#define CUDADataFormats_SiPixelDigi_classes_h
+#ifndef CUDADataFormats_SiPixelDigi_src_classes_h
+#define CUDADataFormats_SiPixelDigi_src_classes_h
 
 #include "CUDADataFormats/Common/interface/Product.h"
-#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "DataFormats/Common/interface/Wrapper.h"
 
-#endif
+#endif  // CUDADataFormats_SiPixelDigi_src_classes_h

CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h

@@ -57,8 +57,10 @@ class TrackSoAT {
 namespace pixelTrack {
 
 #ifdef GPU_SMALL_EVENTS
+  // kept for testing and debugging
   constexpr uint32_t maxNumber() { return 2 * 1024; }
 #else
+  // tested on MC events with 55-75 pileup events
   constexpr uint32_t maxNumber() { return 32 * 1024; }
 #endif
 

CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h

@@ -127,11 +127,7 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(uint32_t nH
   m_hitsLayerStart = view->m_hitsLayerStart = reinterpret_cast<uint32_t*>(get32(n32));
 
   // transfer view
-  if
-#ifndef __CUDACC__
-      constexpr
-#endif
-      (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
+  if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
     cms::cuda::copyAsync(m_view, view, stream);
   } else {
     m_view.reset(view.release());  // NOLINT: std::move() breaks CUDA version

CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h

@@ -14,11 +14,11 @@ namespace pixelCPEforGPU {
 
 class TrackingRecHit2DSOAView {
 public:
-  static constexpr uint32_t maxHits() { return gpuClustering::MaxNumClusters; }
+  static constexpr uint32_t maxHits() { return gpuClustering::maxNumClusters; }
   using hindex_type = uint32_t;  // if above is <=2^32
 
   using PhiBinner =
-      cms::cuda::HistoContainer<int16_t, 128, gpuClustering::MaxNumClusters, 8 * sizeof(int16_t), hindex_type, 10>;
+      cms::cuda::HistoContainer<int16_t, 128, gpuClustering::maxNumClusters, 8 * sizeof(int16_t), hindex_type, 10>;
 
   using Hist = PhiBinner;  // FIXME
 
@@ -89,7 +89,8 @@ class TrackingRecHit2DSOAView {
   uint16_t* m_detInd;
 
   // supporting objects
-  AverageGeometry* m_averageGeometry;  // owned (corrected for beam spot: not sure where to host it otherwise)
+  // m_averageGeometry is corrected for beam spot, not sure where to host it otherwise
+  AverageGeometry* m_averageGeometry;              // owned by TrackingRecHit2DHeterogeneous
   pixelCPEforGPU::ParamsOnGPU const* m_cpeParams;  // forwarded from setup, NOT owned
   uint32_t const* m_hitsModuleStart;               // forwarded from clusters
 

CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc → CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc

@@ -1,4 +1,4 @@
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
@@ -13,7 +13,8 @@ cms::cuda::host::unique_ptr<float[]> TrackingRecHit2DCUDA::localCoordToHostAsync
 
 template <>
 cms::cuda::host::unique_ptr<uint32_t[]> TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cudaStream_t stream) const {
-  auto ret = cms::cuda::make_host_unique<uint32_t[]>(2001, stream);
-  cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, 4 * 2001, cudaMemcpyDefault, stream));
+  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));
   return ret;
 }

CUDADataFormats/TrackingRecHit/src/classes.h

@@ -2,7 +2,7 @@
 #define CUDADataFormats_SiPixelCluster_src_classes_h
 
 #include "CUDADataFormats/Common/interface/Product.h"
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 #include "DataFormats/Common/interface/Wrapper.h"
 
 #endif  // CUDADataFormats_SiPixelCluster_src_classes_h

CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp

@@ -1,4 +1,4 @@
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/requireDevices.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"

CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cu

@@ -1,4 +1,4 @@
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 
 namespace testTrackingRecHit2D {
 

CalibTracker/SiPixelESProducers/plugins/BuildFile.xml

@@ -1,13 +1,15 @@
 <use name="cuda"/>
-<use name="FWCore/Framework"/>
-<use name="FWCore/ParameterSet"/>
+<use name="CalibTracker/Records"/>
+<use name="CalibTracker/SiPixelESProducers"/>
 <use name="CondFormats/DataRecord"/>
 <use name="CondFormats/SiPixelObjects"/>
 <use name="CondFormats/SiStripObjects"/>
+<use name="FWCore/Framework"/>
+<use name="FWCore/ParameterSet"/>
 <use name="Geometry/Records"/>
 <use name="Geometry/TrackerGeometryBuilder"/>
-<use name="CalibTracker/SiPixelESProducers"/>
 <use name="HeterogeneousCore/CUDACore"/>
+<use name="MagneticField/Engine"/>
 <library file="*.cc" name="CalibTrackerSiPixelESProducersPlugins">
   <flags EDM_PLUGIN="1"/>
 </library>