Clean up the pixel local reconstruction code (#593)

Address the pixel local reconstruction review comments. General clean up of the pixel local reconstruction code: - remove commented out and obsolete code and data members - use named constants more consistently - update variable names to follow the coding rules and for better consistency - use member initializer lists in the constructors - allow `if constexpr` in CUDA code - use `std::size` instead of hardcoding the array size - convert iterator-based loops to range-based loops - replace `cout` and `printf` with `LogDebug` or `LogWarning` - use put tokens - reorganise the auto-generated cfi files and use them more consistently - adjust code after rearranging an `#ifdef GPU_DEBUG` block - apply code formatting - other minor changes Improve comments: - improve comments and remove obsolete ones - clarify comments and types regarding `HostProduct` - update comments about `GPU_SMALL_EVENTS` being kept for testing purposes - add notes about the original cpu code Reuse some more common code: - move common pixel cluster code to `PixelClusterizerBase` - extend the `SiPixelCluster` constructor Rename classes and modules for better consistency: - remove the `TrackingRecHit2DCUDA.h` and `gpuClusteringConstants.h` forwarding headers - rename `PixelRecHits` to `PixelRecHitGPUKernel` - rename SiPixelRecHitFromSOA to SiPixelRecHitFromCUDA - rename `siPixelClustersCUDAPreSplitting` to `siPixelClustersPreSplittingCUDA` - rename `siPixelRecHitsCUDAPreSplitting` to `siPixelRecHitsPreSplittingCUDA` - rename `siPixelRecHitsLegacyPreSplitting` to `siPixelRecHitsPreSplittingLegacy` - rename `siPixelRecHitHostSoA` to `siPixelRecHitSoAFromLegacy` Re-apply changes from cms-sw#29805 that were lost in the Patatrack branch.
cms-patatrack · Dec 18, 2020 · 021c154 · 021c154
1 parent 6a192be
commit 021c154
Show file tree

Hide file tree

Showing 79 changed files with 785 additions and 937 deletions.
diff --git a/CUDADataFormats/Common/BuildFile.xml b/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>

diff --git a/CUDADataFormats/SiPixelCluster/BuildFile.xml b/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>
-
diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
@@ -60,4 +60,4 @@ class SiPixelClustersCUDA {
   uint32_t nClusters_h = 0;
 };
 
-#endif
+#endif  // CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/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
 

diff --git a/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc b/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)),

diff --git a/CUDADataFormats/SiPixelCluster/src/classes.h b/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
diff --git a/CUDADataFormats/SiPixelDigi/BuildFile.xml b/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"/>

diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/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
diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h b/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
diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/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)),

diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc b/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);
 }
 

diff --git a/CUDADataFormats/SiPixelDigi/src/classes.h b/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
diff --git a/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h b/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
 

diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h
diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h b/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

diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/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
 

diff --git a/...rackingRecHit/src/TrackingRecHit2DCUDA.cc → ...cHit/src/TrackingRecHit2DHeterogeneous.cc b/...rackingRecHit/src/TrackingRecHit2DCUDA.cc → ...cHit/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;
 }
diff --git a/CUDADataFormats/TrackingRecHit/src/classes.h b/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
diff --git a/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp b/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"

diff --git a/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cu b/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cu
@@ -1,4 +1,4 @@
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 
 namespace testTrackingRecHit2D {
 

diff --git a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml b/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>
-Original file line number
+Diff line change
@@ Expand Up / @@ -60,4 +60,4 @@ class SiPixelClustersCUDA { @@
       uint32_t nClusters_h = 0;
     };
-    #endif
+    #endif  // CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h