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.

SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu

@@ -2,6 +2,7 @@
 #include <limits>
 #include <mutex>
 
+#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudastdAlgorithm.h"
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h"
@@ -19,8 +20,9 @@ __global__ void simLink(const SiPixelDigisCUDA::DeviceConstView* dd,
                         TrackingRecHit2DSOAView const* hhp,
                         ClusterSLView sl,
                         uint32_t n) {
-  constexpr uint32_t invTK = 0;     // std::numeric_limits<int32_t>::max();
-  constexpr uint16_t InvId = 9999;  // must be > MaxNumModules
+  constexpr uint32_t invTK = 0;  // std::numeric_limits<int32_t>::max();
+  using gpuClustering::invalidModuleId;
+  using gpuClustering::maxNumModules;
 
   auto const& hh = *hhp;
   auto i = blockIdx.x * blockDim.x + threadIdx.x;
@@ -29,14 +31,14 @@ __global__ void simLink(const SiPixelDigisCUDA::DeviceConstView* dd,
     return;
 
   auto id = dd->moduleInd(i);
-  if (InvId == id)
+  if (invalidModuleId == id)
     return;
-  assert(id < 2000);
+  assert(id < maxNumModules);
 
   auto ch = pixelgpudetails::pixelToChannel(dd->xx(i), dd->yy(i));
   auto first = hh.hitsModuleStart(id);
   auto cl = first + dd->clus(i);
-  assert(cl < 2000 * blockDim.x);
+  assert(cl < maxNumModules * blockDim.x);
 
   const Clus2TP me{{id, ch, 0, 0, 0, 0, 0}};
 

SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.h

@@ -4,7 +4,7 @@
 #include <cuda_runtime.h>
 
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
 #include "SimTracker/TrackerHitAssociation/interface/trackerHitAssociationHeterogeneous.h"
 

SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc

@@ -6,7 +6,7 @@
 
 #include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
-#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h"
 #include "DataFormats/Common/interface/DetSetVector.h"
 #include "DataFormats/Common/interface/DetSetVectorNew.h"
 #include "DataFormats/Common/interface/Handle.h"
@@ -112,8 +112,8 @@ void ClusterTPAssociationProducerCUDA::fillDescriptions(edm::ConfigurationDescri
   desc.add<edm::InputTag>("stripClusterSrc", edm::InputTag("siStripClusters"));
   desc.add<edm::InputTag>("phase2OTClusterSrc", edm::InputTag("siPhase2Clusters"));
   desc.add<edm::InputTag>("trackingParticleSrc", edm::InputTag("mix", "MergedTrackTruth"));
-  desc.add<edm::InputTag>("heterogeneousPixelDigiClusterSrc", edm::InputTag("siPixelClustersCUDAPreSplitting"));
-  desc.add<edm::InputTag>("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsCUDAPreSplitting"));
+  desc.add<edm::InputTag>("heterogeneousPixelDigiClusterSrc", edm::InputTag("siPixelClustersPreSplittingCUDA"));
+  desc.add<edm::InputTag>("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA"));
 
   desc.add<bool>("dumpCSV", false);