From 0b9b8667ed196543f395c926658c47707adf9bc6 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Fri, 18 Dec 2020 18:26:58 +0100 Subject: [PATCH] Clean up the pixel local reconstruction code (cms-patatrack#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 #29805 that were lost in the Patatrack branch. --- .../TrackerHitAssociation/plugins/ClusterSLOnGPU.cu | 12 +++++++----- .../TrackerHitAssociation/plugins/ClusterSLOnGPU.h | 2 +- .../plugins/ClusterTPAssociationProducerCUDA.cc | 6 +++--- 3 files changed, 11 insertions(+), 9 deletions(-) diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu index 00b03eb86bb34..0aab26d9cc091 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu +++ b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu @@ -2,6 +2,7 @@ #include #include +#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::max(); - constexpr uint16_t InvId = 9999; // must be > MaxNumModules + constexpr uint32_t invTK = 0; // std::numeric_limits::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}}; diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.h b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.h index d8879f2154df4..3109e6ed45a76 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.h +++ b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.h @@ -4,7 +4,7 @@ #include #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" diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc b/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc index cedb9f8fedf29..35337151eda91 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc +++ b/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("stripClusterSrc", edm::InputTag("siStripClusters")); desc.add("phase2OTClusterSrc", edm::InputTag("siPhase2Clusters")); desc.add("trackingParticleSrc", edm::InputTag("mix", "MergedTrackTruth")); - desc.add("heterogeneousPixelDigiClusterSrc", edm::InputTag("siPixelClustersCUDAPreSplitting")); - desc.add("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsCUDAPreSplitting")); + desc.add("heterogeneousPixelDigiClusterSrc", edm::InputTag("siPixelClustersPreSplittingCUDA")); + desc.add("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA")); desc.add("dumpCSV", false);