From f17d75c9b6365c41aab2961ad4d2eb9bb663233a Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 21:35:23 +0100 Subject: [PATCH 01/40] Simple cleanup --- CUDADataFormats/Common/BuildFile.xml | 2 +- CUDADataFormats/SiPixelCluster/BuildFile.xml | 7 +++---- .../SiPixelCluster/interface/SiPixelClustersCUDA.h | 2 +- .../SiPixelCluster/src/SiPixelClustersCUDA.cc | 3 +-- CUDADataFormats/SiPixelCluster/src/classes.h | 6 +++--- CUDADataFormats/SiPixelDigi/BuildFile.xml | 4 ++-- .../SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h | 8 ++++---- .../SiPixelDigi/interface/SiPixelDigisCUDA.h | 6 +++--- .../SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc | 7 +++---- CUDADataFormats/SiPixelDigi/src/classes.h | 8 ++++---- RecoLocalTracker/SiPixelClusterizer/BuildFile.xml | 11 +++++------ .../SiPixelClusterizer/plugins/BuildFile.xml | 8 ++++---- .../SiPixelClusterizer/test/BuildFile.xml | 2 +- RecoLocalTracker/SiPixelRecHits/BuildFile.xml | 12 ++++++------ 14 files changed, 41 insertions(+), 45 deletions(-) diff --git a/CUDADataFormats/Common/BuildFile.xml b/CUDADataFormats/Common/BuildFile.xml index c524cada33060..f6b68fe69b400 100644 --- a/CUDADataFormats/Common/BuildFile.xml +++ b/CUDADataFormats/Common/BuildFile.xml @@ -1,7 +1,7 @@ + - diff --git a/CUDADataFormats/SiPixelCluster/BuildFile.xml b/CUDADataFormats/SiPixelCluster/BuildFile.xml index 5e401d215c4eb..5406d1355533f 100644 --- a/CUDADataFormats/SiPixelCluster/BuildFile.xml +++ b/CUDADataFormats/SiPixelCluster/BuildFile.xml @@ -1,10 +1,9 @@ - - - + + + - diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h index d5d009aaffeb5..acdf1b34a6d79 100644 --- 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/src/SiPixelClustersCUDA.cc b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc index 5e53f49570bb4..ae4a24dbbf83b 100644 --- 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(maxModules + 1, stream)), diff --git a/CUDADataFormats/SiPixelCluster/src/classes.h b/CUDADataFormats/SiPixelCluster/src/classes.h index 0698cb103dab9..3eee5a1fce009 100644 --- 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 index ee357e2d4e157..0806768a9b657 100644 --- a/CUDADataFormats/SiPixelDigi/BuildFile.xml +++ b/CUDADataFormats/SiPixelDigi/BuildFile.xml @@ -1,9 +1,9 @@ + + - - diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h index 85e8883fa1bd4..bfb15c4ac9f5c 100644 --- 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 + #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 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 index 2dc1f628bf426..950f9651cf83b 100644 --- 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 + #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h" -#include - 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 index 70bf2e8aa19f5..eecea35ddd622 100644 --- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc +++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc @@ -1,12 +1,11 @@ -#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h" +#include +#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 - SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream) : data_d(cms::cuda::make_device_unique(maxFedWords, stream)), error_d(cms::cuda::make_device_unique(stream)), diff --git a/CUDADataFormats/SiPixelDigi/src/classes.h b/CUDADataFormats/SiPixelDigi/src/classes.h index fca0811e4650f..fc5d318fad688 100644 --- 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/RecoLocalTracker/SiPixelClusterizer/BuildFile.xml b/RecoLocalTracker/SiPixelClusterizer/BuildFile.xml index 58a7f0b22e30b..7e71c635c95b8 100644 --- a/RecoLocalTracker/SiPixelClusterizer/BuildFile.xml +++ b/RecoLocalTracker/SiPixelClusterizer/BuildFile.xml @@ -1,8 +1,7 @@ - - - - - + + + + - + diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml b/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml index d574c1e6f2b92..a4851e4b322be 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml @@ -1,12 +1,12 @@ - - - - + + + + diff --git a/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml b/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml index 1891970a9d98b..4420adb507027 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml +++ b/RecoLocalTracker/SiPixelClusterizer/test/BuildFile.xml @@ -16,6 +16,7 @@ + @@ -54,4 +55,3 @@ - diff --git a/RecoLocalTracker/SiPixelRecHits/BuildFile.xml b/RecoLocalTracker/SiPixelRecHits/BuildFile.xml index e22b18b17117a..d9376d88f7bbd 100644 --- a/RecoLocalTracker/SiPixelRecHits/BuildFile.xml +++ b/RecoLocalTracker/SiPixelRecHits/BuildFile.xml @@ -1,14 +1,14 @@ - - + + + + + - + - - - From 115a591cc27a9b1fce842368ff40bf6555b601fe Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 22:10:57 +0100 Subject: [PATCH 02/40] Rename gpuClusteringConstants to lowercase --- .../interface/gpuClusteringConstants.h | 11 ++-- .../interface/TrackingRecHit2DSOAView.h | 4 +- .../plugins/SiPixelRawToClusterGPUKernel.cu | 31 ++++++----- .../plugins/gpuCalibPixel.h | 9 ++-- .../plugins/gpuClusterChargeCut.h | 36 ++++++------- .../plugins/gpuClustering.h | 22 ++++---- .../SiPixelClusterizer/test/gpuClustering_t.h | 54 +++++++++---------- .../plugins/SiPixelRecHitConverter.cc | 10 ++-- .../plugins/SiPixelRecHitFromSOA.cc | 10 ++-- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 14 ++--- .../SiPixelRecHits/plugins/gpuPixelRecHits.h | 9 ++-- .../plugins/PixelTrackProducerFromSoA.cc | 2 +- .../plugins/ClusterSLOnGPU.cu | 7 +-- 13 files changed, 111 insertions(+), 108 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index 1430606ab6678..5928d45af7dc8 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -2,6 +2,7 @@ #define CUDADataFormats_SiPixelCluster_interface_gpuClusteringConstants_h #include +#include namespace pixelGPUConstants { #ifdef GPU_SMALL_EVENTS @@ -21,11 +22,11 @@ namespace gpuClustering { #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::max() - 1; + static_assert(invalidModuleId > maxNumModules); // invalidModuleId must be > maxNumModules } // namespace gpuClustering diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h index 6a83a66b60fbd..6335a4f9346bf 100644 --- 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; + cms::cuda::HistoContainer; using Hist = PhiBinner; // FIXME diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 7c133b10f4dab..b1ab47be7dad6 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -472,22 +472,22 @@ namespace pixelgpudetails { } // end of Raw to Digi kernel __global__ void fillHitsModuleStart(uint32_t const *__restrict__ cluStart, uint32_t *__restrict__ moduleStart) { - assert(gpuClustering::MaxNumModules < 2048); // easy to extend at least till 32*1024 + assert(gpuClustering::maxNumModules < 2048); // easy to extend at least till 32*1024 assert(1 == gridDim.x); assert(0 == blockIdx.x); int first = threadIdx.x; - // limit to MaxHitsInModule; - for (int i = first, iend = gpuClustering::MaxNumModules; i < iend; i += blockDim.x) { + // limit to maxHitsInModule() + for (int i = first, iend = gpuClustering::maxNumModules; i < iend; i += blockDim.x) { moduleStart[i + 1] = std::min(gpuClustering::maxHitsInModule(), cluStart[i]); } __shared__ uint32_t ws[32]; cms::cuda::blockPrefixScan(moduleStart + 1, moduleStart + 1, 1024, ws); - cms::cuda::blockPrefixScan(moduleStart + 1025, moduleStart + 1025, gpuClustering::MaxNumModules - 1024, ws); + cms::cuda::blockPrefixScan(moduleStart + 1025, moduleStart + 1025, gpuClustering::maxNumModules - 1024, ws); - for (int i = first + 1025, iend = gpuClustering::MaxNumModules + 1; i < iend; i += blockDim.x) { + for (int i = first + 1025, iend = gpuClustering::maxNumModules + 1; i < iend; i += blockDim.x) { moduleStart[i] += moduleStart[1024]; } __syncthreads(); @@ -498,23 +498,22 @@ namespace pixelgpudetails { assert(c0 == moduleStart[1]); assert(moduleStart[1024] >= moduleStart[1023]); assert(moduleStart[1025] >= moduleStart[1024]); - assert(moduleStart[gpuClustering::MaxNumModules] >= moduleStart[1025]); + assert(moduleStart[gpuClustering::maxNumModules] >= moduleStart[1025]); - for (int i = first, iend = gpuClustering::MaxNumModules + 1; i < iend; i += blockDim.x) { + for (int i = first, iend = gpuClustering::maxNumModules + 1; i < iend; i += blockDim.x) { if (0 != i) assert(moduleStart[i] >= moduleStart[i - i]); // [BPX1, BPX2, BPX3, BPX4, FP1, FP2, FP3, FN1, FN2, FN3, LAST_VALID] // [ 0, 96, 320, 672, 1184, 1296, 1408, 1520, 1632, 1744, 1856] - if (i == 96 || i == 1184 || i == 1744 || i == gpuClustering::MaxNumModules) + if (i == 96 || i == 1184 || i == 1744 || i == gpuClustering::maxNumModules) printf("moduleStart %d %d\n", i, moduleStart[i]); } #endif // avoid overflow - constexpr auto MAX_HITS = gpuClustering::MaxNumClusters; - for (int i = first, iend = gpuClustering::MaxNumModules + 1; i < iend; i += blockDim.x) { - if (moduleStart[i] > MAX_HITS) - moduleStart[i] = MAX_HITS; + auto constexpr maxNumClusters = gpuClustering::maxNumClusters; + for (int i = first, iend = gpuClustering::maxNumModules + 1; i < iend; i += blockDim.x) { + moduleStart[i] = std::clamp(moduleStart[i], 0U, maxNumClusters); } } @@ -541,7 +540,7 @@ namespace pixelgpudetails { if (includeErrors) { digiErrors_d = SiPixelDigiErrorsCUDA(pixelgpudetails::MAX_FED_WORDS, std::move(errors), stream); } - clusters_d = SiPixelClustersCUDA(gpuClustering::MaxNumModules, stream); + clusters_d = SiPixelClustersCUDA(gpuClustering::maxNumModules, stream); nModules_Clusters_h = cms::cuda::make_host_unique(2, stream); @@ -594,7 +593,7 @@ namespace pixelgpudetails { using namespace gpuClustering; int threadsPerBlock = 256; int blocks = - (std::max(int(wordCounter), int(gpuClustering::MaxNumModules)) + threadsPerBlock - 1) / threadsPerBlock; + (std::max(int(wordCounter), int(gpuClustering::maxNumModules)) + threadsPerBlock - 1) / threadsPerBlock; gpuCalibPixel::calibDigis<<>>(isRun2, digis_d.moduleInd(), @@ -626,7 +625,7 @@ namespace pixelgpudetails { &(nModules_Clusters_h[0]), clusters_d.moduleStart(), sizeof(uint32_t), cudaMemcpyDefault, stream)); threadsPerBlock = 256; - blocks = MaxNumModules; + blocks = maxNumModules; #ifdef GPU_DEBUG std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif @@ -664,7 +663,7 @@ namespace pixelgpudetails { // last element holds the number of all clusters cudaCheck(cudaMemcpyAsync(&(nModules_Clusters_h[1]), - clusters_d.clusModuleStart() + gpuClustering::MaxNumModules, + clusters_d.clusModuleStart() + gpuClustering::maxNumModules, sizeof(uint32_t), cudaMemcpyDefault, stream)); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h index 50c62f44f1df8..572a482144667 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h @@ -4,6 +4,7 @@ #include #include +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" @@ -11,7 +12,7 @@ namespace gpuCalibPixel { - constexpr uint16_t InvId = 9999; // must be > MaxNumModules + using gpuClustering::invalidModuleId; // valid for run2 constexpr float VCaltoElectronGain = 47; // L2-4: 47 +- 4.7 @@ -35,12 +36,12 @@ namespace gpuCalibPixel { // zero for next kernels... if (0 == first) clusModuleStart[0] = moduleStart[0] = 0; - for (int i = first; i < gpuClustering::MaxNumModules; i += gridDim.x * blockDim.x) { + for (int i = first; i < gpuClustering::maxNumModules; i += gridDim.x * blockDim.x) { nClustersInModule[i] = 0; } for (int i = first; i < numElements; i += gridDim.x * blockDim.x) { - if (InvId == id[i]) + if (invalidModuleId == id[i]) continue; float conversionFactor = (isRun2) ? (id[i] < 96 ? VCaltoElectronGain_L1 : VCaltoElectronGain) : 1.f; @@ -55,7 +56,7 @@ namespace gpuCalibPixel { float gain = ret.second; // float pedestal = 0; float gain = 1.; if (isDeadColumn | isNoisyColumn) { - id[i] = InvId; + id[i] = invalidModuleId; adc[i] = 0; printf("bad pixel at %d in %d\n", i, id[i]); } else { diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h index 8f45d35b267b1..8a90134488ec5 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -19,44 +19,44 @@ namespace gpuClustering { uint32_t const* __restrict__ moduleId, // module id of each module int32_t* __restrict__ clusterId, // modified: cluster id of each pixel uint32_t numElements) { - __shared__ int32_t charge[MaxNumClustersPerModules]; - __shared__ uint8_t ok[MaxNumClustersPerModules]; - __shared__ uint16_t newclusId[MaxNumClustersPerModules]; + __shared__ int32_t charge[maxNumClustersPerModules]; + __shared__ uint8_t ok[maxNumClustersPerModules]; + __shared__ uint16_t newclusId[maxNumClustersPerModules]; auto firstModule = blockIdx.x; auto endModule = moduleStart[0]; for (auto module = firstModule; module < endModule; module += gridDim.x) { auto firstPixel = moduleStart[1 + module]; auto thisModuleId = id[firstPixel]; - assert(thisModuleId < MaxNumModules); + assert(thisModuleId < maxNumModules); assert(thisModuleId == moduleId[module]); auto nclus = nClustersInModule[thisModuleId]; if (nclus == 0) continue; - if (threadIdx.x == 0 && nclus > MaxNumClustersPerModules) + if (threadIdx.x == 0 && nclus > maxNumClustersPerModules) printf("Warning too many clusters in module %d in block %d: %d > %d\n", thisModuleId, blockIdx.x, nclus, - MaxNumClustersPerModules); + maxNumClustersPerModules); auto first = firstPixel + threadIdx.x; - if (nclus > MaxNumClustersPerModules) { + if (nclus > maxNumClustersPerModules) { // remove excess FIXME find a way to cut charge first.... for (auto i = first; i < numElements; i += blockDim.x) { - if (id[i] == InvId) + if (id[i] == invalidModuleId) continue; // not valid if (id[i] != thisModuleId) break; // end of module - if (clusterId[i] >= MaxNumClustersPerModules) { - id[i] = InvId; - clusterId[i] = InvId; + if (clusterId[i] >= maxNumClustersPerModules) { + id[i] = invalidModuleId; + clusterId[i] = invalidModuleId; } } - nclus = MaxNumClustersPerModules; + nclus = maxNumClustersPerModules; } #ifdef GPU_DEBUG @@ -65,14 +65,14 @@ namespace gpuClustering { printf("start cluster charge cut for module %d in block %d\n", thisModuleId, blockIdx.x); #endif - assert(nclus <= MaxNumClustersPerModules); + assert(nclus <= maxNumClustersPerModules); for (auto i = threadIdx.x; i < nclus; i += blockDim.x) { charge[i] = 0; } __syncthreads(); for (auto i = first; i < numElements; i += blockDim.x) { - if (id[i] == InvId) + if (id[i] == invalidModuleId) continue; // not valid if (id[i] != thisModuleId) break; // end of module @@ -102,19 +102,19 @@ namespace gpuClustering { // mark bad cluster again for (auto i = threadIdx.x; i < nclus; i += blockDim.x) { if (0 == ok[i]) - newclusId[i] = InvId + 1; + newclusId[i] = invalidModuleId + 1; } __syncthreads(); // reassign id for (auto i = first; i < numElements; i += blockDim.x) { - if (id[i] == InvId) + if (id[i] == invalidModuleId) continue; // not valid if (id[i] != thisModuleId) break; // end of module clusterId[i] = newclusId[clusterId[i]] - 1; - if (clusterId[i] == InvId) - id[i] = InvId; + if (clusterId[i] == invalidModuleId) + id[i] = invalidModuleId; } //done diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h index 3bf42c8265b1e..508138f5fbb57 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h @@ -23,14 +23,14 @@ namespace gpuClustering { int first = blockDim.x * blockIdx.x + threadIdx.x; for (int i = first; i < numElements; i += gridDim.x * blockDim.x) { clusterId[i] = i; - if (InvId == id[i]) + if (invalidModuleId == id[i]) continue; auto j = i - 1; - while (j >= 0 and id[j] == InvId) + while (j >= 0 and id[j] == invalidModuleId) --j; if (j < 0 or id[j] != id[i]) { // boundary... - auto loc = atomicInc(moduleStart, MaxNumModules); + auto loc = atomicInc(moduleStart, maxNumModules); moduleStart[loc + 1] = i; } } @@ -54,7 +54,7 @@ namespace gpuClustering { for (auto module = firstModule; module < endModule; module += gridDim.x) { auto firstPixel = moduleStart[1 + module]; auto thisModuleId = id[firstPixel]; - assert(thisModuleId < MaxNumModules); + assert(thisModuleId < maxNumModules); #ifdef GPU_DEBUG if (thisModuleId % 100 == 1) @@ -70,7 +70,7 @@ namespace gpuClustering { // skip threads not associated to an existing pixel for (int i = first; i < numElements; i += blockDim.x) { - if (id[i] == InvId) // skip invalid pixels + if (id[i] == invalidModuleId) // skip invalid pixels continue; if (id[i] != thisModuleId) { // find the first pixel in a different module atomicMin(&msize, i); @@ -110,7 +110,7 @@ namespace gpuClustering { // fill histo for (int i = first; i < msize; i += blockDim.x) { - if (id[i] == InvId) // skip invalid pixels + if (id[i] == invalidModuleId) // skip invalid pixels continue; hist.count(y[i]); #ifdef GPU_DEBUG @@ -130,7 +130,7 @@ namespace gpuClustering { printf("histo size %d\n", hist.size()); #endif for (int i = first; i < msize; i += blockDim.x) { - if (id[i] == InvId) // skip invalid pixels + if (id[i] == invalidModuleId) // skip invalid pixels continue; hist.fill(y[i], i - firstPixel); } @@ -178,7 +178,7 @@ namespace gpuClustering { assert(k < maxiter); auto p = hist.begin() + j; auto i = *p + firstPixel; - assert(id[i] != InvId); + assert(id[i] != invalidModuleId); assert(id[i] == thisModuleId); // same module int be = Hist::bin(y[i] + 1); auto e = hist.end(be); @@ -255,7 +255,7 @@ namespace gpuClustering { // find the number of different clusters, identified by a pixels with clus[i] == i; // mark these pixels with a negative id. for (int i = first; i < msize; i += blockDim.x) { - if (id[i] == InvId) // skip invalid pixels + if (id[i] == invalidModuleId) // skip invalid pixels continue; if (clusterId[i] == i) { auto old = atomicInc(&foundClusters, 0xffffffff); @@ -266,7 +266,7 @@ namespace gpuClustering { // propagate the negative id to all the pixels in the cluster. for (int i = first; i < msize; i += blockDim.x) { - if (id[i] == InvId) // skip invalid pixels + if (id[i] == invalidModuleId) // skip invalid pixels continue; if (clusterId[i] >= 0) { // mark each pixel in a cluster with the same id as the first one @@ -277,7 +277,7 @@ namespace gpuClustering { // adjust the cluster id to be a positive value starting from 0 for (int i = first; i < msize; i += blockDim.x) { - if (id[i] == InvId) { // skip invalid pixels + if (id[i] == invalidModuleId) { // skip invalid pixels clusterId[i] = -9999; continue; } diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h index 64289d5208b48..e3e5f17604df0 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h +++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h @@ -40,13 +40,13 @@ int main(void) { auto d_y = cms::cuda::make_device_unique(numElements, nullptr); auto d_adc = cms::cuda::make_device_unique(numElements, nullptr); auto d_clus = cms::cuda::make_device_unique(numElements, nullptr); - auto d_moduleStart = cms::cuda::make_device_unique(MaxNumModules + 1, nullptr); - auto d_clusInModule = cms::cuda::make_device_unique(MaxNumModules, nullptr); - auto d_moduleId = cms::cuda::make_device_unique(MaxNumModules, nullptr); -#else // __CUDACC__ - auto h_moduleStart = std::make_unique(MaxNumModules + 1); - auto h_clusInModule = std::make_unique(MaxNumModules); - auto h_moduleId = std::make_unique(MaxNumModules); + auto d_moduleStart = cms::cuda::make_device_unique(maxNumModules + 1, nullptr); + auto d_clusInModule = cms::cuda::make_device_unique(maxNumModules, nullptr); + auto d_moduleId = cms::cuda::make_device_unique(maxNumModules, nullptr); +#else // __CUDACC__ + auto h_moduleStart = std::make_unique(maxNumModules + 1); + auto h_clusInModule = std::make_unique(maxNumModules); + auto h_moduleId = std::make_unique(maxNumModules); #endif // __CUDACC__ // later random number @@ -145,7 +145,7 @@ int main(void) { ++n; } ++ncl; - h_id[n++] = InvId; // error + h_id[n++] = invalidModuleId; // error // messy int xx[5] = {21, 25, 23, 24, 22}; for (int k = 0; k < 5; ++k) { @@ -186,7 +186,7 @@ int main(void) { // all odd id for (int id = 11; id <= 1800; id += 2) { if ((id / 20) % 2) - h_id[n++] = InvId; // error + h_id[n++] = invalidModuleId; // error for (int x = 0; x < 40; x += 4) { ++ncl; if ((id / 10) % 2) { @@ -212,8 +212,8 @@ int main(void) { if (y[k] == 3) continue; // hole if (id == 51) { - h_id[n++] = InvId; - h_id[n++] = InvId; + h_id[n++] = invalidModuleId; + h_id[n++] = invalidModuleId; } // error h_id[n] = id; h_x[n] = x + 1; @@ -253,11 +253,11 @@ int main(void) { cms::cuda::launch(countModules, {blocksPerGrid, threadsPerBlock}, d_id.get(), d_moduleStart.get(), d_clus.get(), n); - blocksPerGrid = MaxNumModules; //nModules; + blocksPerGrid = maxNumModules; //nModules; std::cout << "CUDA findModules kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads\n"; - cudaCheck(cudaMemset(d_clusInModule.get(), 0, MaxNumModules * sizeof(uint32_t))); + cudaCheck(cudaMemset(d_clusInModule.get(), 0, maxNumModules * sizeof(uint32_t))); cms::cuda::launch(findClus, {blocksPerGrid, threadsPerBlock}, @@ -272,17 +272,17 @@ int main(void) { cudaDeviceSynchronize(); cudaCheck(cudaMemcpy(&nModules, d_moduleStart.get(), sizeof(uint32_t), cudaMemcpyDeviceToHost)); - uint32_t nclus[MaxNumModules], moduleId[nModules]; - cudaCheck(cudaMemcpy(&nclus, d_clusInModule.get(), MaxNumModules * sizeof(uint32_t), cudaMemcpyDeviceToHost)); + uint32_t nclus[maxNumModules], moduleId[nModules]; + cudaCheck(cudaMemcpy(&nclus, d_clusInModule.get(), maxNumModules * sizeof(uint32_t), cudaMemcpyDeviceToHost)); - std::cout << "before charge cut found " << std::accumulate(nclus, nclus + MaxNumModules, 0) << " clusters" + std::cout << "before charge cut found " << std::accumulate(nclus, nclus + maxNumModules, 0) << " clusters" << std::endl; - for (auto i = MaxNumModules; i > 0; i--) + for (auto i = maxNumModules; i > 0; i--) if (nclus[i - 1] > 0) { std::cout << "last module is " << i - 1 << ' ' << nclus[i - 1] << std::endl; break; } - if (ncl != std::accumulate(nclus, nclus + MaxNumModules, 0)) + if (ncl != std::accumulate(nclus, nclus + maxNumModules, 0)) std::cout << "ERROR!!!!! wrong number of cluster found" << std::endl; cms::cuda::launch(clusterChargeCut, @@ -296,24 +296,24 @@ int main(void) { n); cudaDeviceSynchronize(); -#else // __CUDACC__ +#else // __CUDACC__ h_moduleStart[0] = nModules; countModules(h_id.get(), h_moduleStart.get(), h_clus.get(), n); - memset(h_clusInModule.get(), 0, MaxNumModules * sizeof(uint32_t)); + memset(h_clusInModule.get(), 0, maxNumModules * sizeof(uint32_t)); findClus( h_id.get(), h_x.get(), h_y.get(), h_moduleStart.get(), h_clusInModule.get(), h_moduleId.get(), h_clus.get(), n); nModules = h_moduleStart[0]; auto nclus = h_clusInModule.get(); - std::cout << "before charge cut found " << std::accumulate(nclus, nclus + MaxNumModules, 0) << " clusters" + std::cout << "before charge cut found " << std::accumulate(nclus, nclus + maxNumModules, 0) << " clusters" << std::endl; - for (auto i = MaxNumModules; i > 0; i--) + for (auto i = maxNumModules; i > 0; i--) if (nclus[i - 1] > 0) { std::cout << "last module is " << i - 1 << ' ' << nclus[i - 1] << std::endl; break; } - if (ncl != std::accumulate(nclus, nclus + MaxNumModules, 0)) + if (ncl != std::accumulate(nclus, nclus + maxNumModules, 0)) std::cout << "ERROR!!!!! wrong number of cluster found" << std::endl; clusterChargeCut( @@ -325,14 +325,14 @@ int main(void) { #ifdef __CUDACC__ cudaCheck(cudaMemcpy(h_id.get(), d_id.get(), size16, cudaMemcpyDeviceToHost)); cudaCheck(cudaMemcpy(h_clus.get(), d_clus.get(), size32, cudaMemcpyDeviceToHost)); - cudaCheck(cudaMemcpy(&nclus, d_clusInModule.get(), MaxNumModules * sizeof(uint32_t), cudaMemcpyDeviceToHost)); + cudaCheck(cudaMemcpy(&nclus, d_clusInModule.get(), maxNumModules * sizeof(uint32_t), cudaMemcpyDeviceToHost)); cudaCheck(cudaMemcpy(&moduleId, d_moduleId.get(), nModules * sizeof(uint32_t), cudaMemcpyDeviceToHost)); #endif // __CUDACC__ std::set clids; for (int i = 0; i < n; ++i) { assert(h_id[i] != 666); // only noise - if (h_id[i] == InvId) + if (h_id[i] == invalidModuleId) continue; assert(h_clus[i] >= 0); assert(h_clus[i] < int(nclus[h_id[i]])); @@ -368,9 +368,9 @@ int main(void) { std::cout << "error " << mid << ": " << nc << ' ' << pnc << std::endl; } - std::cout << "found " << std::accumulate(nclus, nclus + MaxNumModules, 0) << ' ' << clids.size() << " clusters" + std::cout << "found " << std::accumulate(nclus, nclus + maxNumModules, 0) << ' ' << clids.size() << " clusters" << std::endl; - for (auto i = MaxNumModules; i > 0; i--) + for (auto i = maxNumModules; i > 0; i--) if (nclus[i - 1] > 0) { std::cout << "last module is " << i - 1 << ' ' << nclus[i - 1] << std::endl; break; diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc index c7eb7481fc4f8..4ebba652335a0 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc @@ -190,16 +190,16 @@ namespace cms { const edmNew::DetSetVector& input = *inputhandle; // fill cluster arrays - auto hmsp = std::make_unique(gpuClustering::MaxNumModules + 1); + auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); auto hitsModuleStart = hmsp.get(); - std::array clusInModule{}; + std::array clusInModule{}; for (auto DSViter = input.begin(); DSViter != input.end(); DSViter++) { unsigned int detid = DSViter->detId(); DetId detIdObject(detid); const GeomDetUnit* genericDet = geom.idToDetUnit(detIdObject); auto gind = genericDet->index(); // FIXME to be changed to support Phase2 - if (gind >= int(gpuClustering::MaxNumModules)) + if (gind >= int(gpuClustering::maxNumModules)) continue; auto const nclus = DSViter->size(); assert(nclus > 0); @@ -207,10 +207,10 @@ namespace cms { numberOfClusters += nclus; } hitsModuleStart[0] = 0; - assert(clusInModule.size() > gpuClustering::MaxNumModules); + assert(clusInModule.size() > gpuClustering::maxNumModules); for (int i = 1, n = clusInModule.size(); i < n; ++i) hitsModuleStart[i] = hitsModuleStart[i - 1] + clusInModule[i - 1]; - assert(numberOfClusters == int(hitsModuleStart[gpuClustering::MaxNumModules])); + assert(numberOfClusters == int(hitsModuleStart[gpuClustering::maxNumModules])); // yes a unique ptr of a unique ptr so edm is happy and the pointer stay still... iEvent.emplace(tHost_, std::move(hmsp)); // hmsp is gone, hitsModuleStart still alive and kicking... diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc index f9e9825aa7be8..a596c1030619e 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc @@ -83,7 +83,7 @@ void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es) { // yes a unique ptr of a unique ptr so edm is happy - auto sizeOfHitModuleStart = gpuClustering::MaxNumModules + 1; + auto sizeOfHitModuleStart = gpuClustering::maxNumModules + 1; auto hmsp = std::make_unique(sizeOfHitModuleStart); std::copy(m_hitsModuleStart.get(), m_hitsModuleStart.get() + sizeOfHitModuleStart, hmsp.get()); auto hms = std::make_unique(std::move(hmsp)); // hmsp is gone @@ -107,7 +107,7 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es auto const& input = *hclusters; - constexpr uint32_t MaxHitsInModule = gpuClustering::MaxHitsInModule; + constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule(); int numberOfDetUnits = 0; int numberOfClusters = 0; @@ -127,10 +127,10 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es assert(lc > fc); // std::cout << "in det " << gind << ": conv " << nhits << " hits from " << DSViter->size() << " legacy clusters" // <<' '<< fc <<','< MaxHitsInModule) + if (nhits > maxHitsInModule) printf( - "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nhits, gind, MaxHitsInModule); - nhits = std::min(nhits, MaxHitsInModule); + "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nhits, gind, maxHitsInModule); + nhits = std::min(nhits, maxHitsInModule); //std::cout << "in det " << gind << "conv " << nhits << " hits from " << DSViter->size() << " legacy clusters" // <<' '<< lc <<','< { static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); - using HitModuleStart = std::array; + using HitModuleStart = std::array; using HMSstorage = HostProduct; private: @@ -91,7 +91,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv auto const& input = *hclusters; // yes a unique ptr of a unique ptr so edm is happy and the pointer stay still... - auto hmsp = std::make_unique(gpuClustering::MaxNumModules + 1); + auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); auto hitsModuleStart = hmsp.get(); auto hms = std::make_unique(std::move(hmsp)); // hmsp is gone iEvent.put(tokenModuleStart_, std::move(hms)); // hms is gone! hitsModuleStart still alive and kicking... @@ -108,7 +108,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv std::vector, SiPixelCluster>> clusterRef; - constexpr uint32_t MaxHitsInModule = gpuClustering::MaxHitsInModule; + constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule(); HitModuleStart moduleStart_; // index of the first pixel of each module HitModuleStart clusInModule_; @@ -173,9 +173,9 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv // std::cout << "in det " << gind << ": conv " << nclus << " hits from " << DSViter->size() << " legacy clusters" // <<' '<< fc <<','< MaxHitsInModule) + if (nclus > maxHitsInModule) printf( - "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nclus, gind, MaxHitsInModule); + "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nclus, gind, maxHitsInModule); // fill digis xx_.clear(); @@ -212,7 +212,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv // we run on blockId.x==0 gpuPixelRecHits::getHits(&cpeView, &bsHost, &digiView, ndigi, &clusterView, output->view()); for (auto h = fc; h < lc; ++h) - if (h - fc < MaxHitsInModule) + if (h - fc < maxHitsInModule) assert(gind == output->view()->detectorIndex(h)); else assert(9999 == output->view()->detectorIndex(h)); @@ -220,7 +220,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv SiPixelRecHitCollectionNew::FastFiller recHitsOnDetUnit(*legacyOutput, detid); for (auto h = fc; h < lc; ++h) { auto ih = h - fc; - if (ih >= MaxHitsInModule) + if (ih >= maxHitsInModule) break; assert(ih < clusterRef.size()); LocalPoint lp(output->view()->xLocal(h), output->view()->yLocal(h)); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index 17cd5aad4db52..eddde2237f399 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -6,6 +6,7 @@ #include #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" #include "DataFormats/Math/interface/approx_atan2.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" @@ -53,7 +54,7 @@ namespace gpuPixelRecHits { } // to be moved in common namespace... - constexpr uint16_t InvId = 9999; // must be > MaxNumModules + using gpuClustering::invalidModuleId; constexpr int32_t MaxHitsInIter = pixelCPEforGPU::MaxHitsInIter; using ClusParams = pixelCPEforGPU::ClusParams; @@ -70,7 +71,7 @@ namespace gpuPixelRecHits { #ifdef GPU_DEBUG if (threadIdx.x == 0) { auto k = first; - while (digis.moduleInd(k) == InvId) + while (digis.moduleInd(k) == invalidModuleId) ++k; assert(digis.moduleInd(k) == me); } @@ -114,7 +115,7 @@ namespace gpuPixelRecHits { for (int i = first; i < numElements; i += blockDim.x) { auto id = digis.moduleInd(i); - if (id == InvId) + if (id == invalidModuleId) continue; // not valid if (id != me) break; // end of module @@ -137,7 +138,7 @@ namespace gpuPixelRecHits { auto pixmx = cpeParams->detParams(me).pixmx; for (int i = first; i < numElements; i += blockDim.x) { auto id = digis.moduleInd(i); - if (id == InvId) + if (id == invalidModuleId) continue; // not valid if (id != me) break; // end of module diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index 522678ce352f5..005651ab14493 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc @@ -47,7 +47,7 @@ class PixelTrackProducerFromSoA : public edm::global::EDProducer<> { static void fillDescriptions(edm::ConfigurationDescriptions &descriptions); - // using HitModuleStart = std::array; + // using HitModuleStart = std::array; using HMSstorage = HostProduct; private: diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu index 00b03eb86bb34..c06d6d254bad3 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,8 @@ __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; auto const& hh = *hhp; auto i = blockIdx.x * blockDim.x + threadIdx.x; @@ -29,7 +30,7 @@ __global__ void simLink(const SiPixelDigisCUDA::DeviceConstView* dd, return; auto id = dd->moduleInd(i); - if (InvId == id) + if (invalidModuleId == id) return; assert(id < 2000); From c962e7932e0ea955daecb929004cfee598c43702 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 22:29:00 +0100 Subject: [PATCH 03/40] Use member initializer list in the constructor --- .../SiPixelDigi/src/SiPixelDigisCUDA.cc | 24 ++++++++----------- 1 file changed, 10 insertions(+), 14 deletions(-) diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc index 664364b6ff25a..4e6a3fc2593fd 100644 --- 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(maxFedWords, stream); - yy_d = cms::cuda::make_device_unique(maxFedWords, stream); - adc_d = cms::cuda::make_device_unique(maxFedWords, stream); - moduleInd_d = cms::cuda::make_device_unique(maxFedWords, stream); - clus_d = cms::cuda::make_device_unique(maxFedWords, stream); - - pdigi_d = cms::cuda::make_device_unique(maxFedWords, stream); - rawIdArr_d = cms::cuda::make_device_unique(maxFedWords, stream); +SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream) + : xx_d(cms::cuda::make_device_unique(maxFedWords, stream)), + yy_d(cms::cuda::make_device_unique(maxFedWords, stream)), + adc_d(cms::cuda::make_device_unique(maxFedWords, stream)), + moduleInd_d(cms::cuda::make_device_unique(maxFedWords, stream)), + clus_d(cms::cuda::make_device_unique(maxFedWords, stream)), + view_d(cms::cuda::make_device_unique(stream)), + pdigi_d(cms::cuda::make_device_unique(maxFedWords, stream)), + rawIdArr_d(cms::cuda::make_device_unique(maxFedWords, stream)) { auto view = cms::cuda::make_host_unique(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(stream); cms::cuda::copyAsync(view_d, view, stream); } From 231640127c11decf53312a84678bd99497d0c683 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 22:37:08 +0100 Subject: [PATCH 04/40] Drop TrackingRecHit2DCUDA.h compatibility header --- CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h | 1 - ...TrackingRecHit2DCUDA.cc => TrackingRecHit2DHeterogeneous.cc} | 2 +- CUDADataFormats/TrackingRecHit/src/classes.h | 2 +- CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp | 2 +- CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cu | 2 +- RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h | 2 +- RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc | 2 +- RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc | 2 +- RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h | 2 +- .../PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc | 2 +- RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h | 2 +- RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc | 2 +- .../PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h | 2 +- RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h | 2 +- RecoPixelVertexing/PixelTriplets/plugins/HelixFitOnGPU.h | 2 +- RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h | 2 +- .../PixelTriplets/plugins/gpuPixelDoubletsAlgos.h | 2 +- SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.h | 2 +- .../plugins/ClusterTPAssociationProducerCUDA.cc | 2 +- 19 files changed, 18 insertions(+), 19 deletions(-) delete mode 100644 CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h rename CUDADataFormats/TrackingRecHit/src/{TrackingRecHit2DCUDA.cc => TrackingRecHit2DHeterogeneous.cc} (97%) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h deleted file mode 100644 index f6b715b3e743e..0000000000000 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h +++ /dev/null @@ -1 +0,0 @@ -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" diff --git a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc similarity index 97% rename from CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc rename to CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc index 7b04ed2d530a0..d4bf1b500e216 100644 --- a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc +++ b/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" diff --git a/CUDADataFormats/TrackingRecHit/src/classes.h b/CUDADataFormats/TrackingRecHit/src/classes.h index 3d40821493c5b..86fef25746efd 100644 --- 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 index 32af6c181ae68..3d8413b36ec96 100644 --- 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 index 6b55f8a8f98c5..06bd599d074f9 100644 --- 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/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h index 594057432d9a2..ac02caab3259e 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h @@ -8,7 +8,7 @@ #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" namespace pixelgpudetails { diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc index 10e4b678b515c..1312126be32c1 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc @@ -4,7 +4,7 @@ #include "CUDADataFormats/Common/interface/Product.h" #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "DataFormats/Common/interface/Handle.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/EventSetup.h" diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc index a596c1030619e..b2e7c36cffd10 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc @@ -2,7 +2,7 @@ #include "CUDADataFormats/Common/interface/Product.h" #include "CUDADataFormats/Common/interface/HostProduct.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "DataFormats/Common/interface/DetSetVectorNew.h" #include "DataFormats/Common/interface/Handle.h" #include "DataFormats/SiPixelCluster/interface/SiPixelCluster.h" diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index eddde2237f399..1db62e6ead85a 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -7,7 +7,7 @@ #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "DataFormats/Math/interface/approx_atan2.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h" diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc index 04faf570c3fcc..2f0965be50eb8 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackDumpCUDA.cc @@ -2,7 +2,7 @@ #include "CUDADataFormats/Common/interface/Product.h" #include "CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "CUDADataFormats/Vertex/interface/ZVertexHeterogeneous.h" #include "DataFormats/Common/interface/Handle.h" #include "FWCore/Framework/interface/ConsumesCollector.h" diff --git a/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h index 7b5c9ea0ce0a4..96a641829d797 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/BrokenLineFitOnGPU.h @@ -8,7 +8,7 @@ #include -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h" diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc index 31e5070e55e05..3f6ea5f43c6f9 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc @@ -20,7 +20,7 @@ #include "CAHitNtupletGeneratorOnGPU.h" #include "CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" class CAHitNtupletCUDA : public edm::global::EDProducer<> { public: diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h index e920ebf7a803d..afb591744bf59 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.h @@ -2,7 +2,7 @@ #define RecoPixelVertexing_PixelTriplets_plugins_CAHitNtupletGeneratorOnGPU_h #include -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h" #include "DataFormats/SiPixelDetId/interface/PixelSubdetector.h" diff --git a/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h b/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h index e913b77fe0953..ef600489f3e0f 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h @@ -9,7 +9,7 @@ #include -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h" #include "HeterogeneousCore/CUDAUtilities/interface/VecArray.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" diff --git a/RecoPixelVertexing/PixelTriplets/plugins/HelixFitOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/HelixFitOnGPU.h index 05b399e870f58..42f8f0e720b43 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/HelixFitOnGPU.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/HelixFitOnGPU.h @@ -2,7 +2,7 @@ #define RecoPixelVertexing_PixelTrackFitting_plugins_HelixFitOnGPU_h #include "CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h" -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "RecoPixelVertexing/PixelTrackFitting/interface/FitResult.h" #include "CAConstants.h" diff --git a/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h b/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h index 75e9d570a129e..a16374278233a 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/RiemannFitOnGPU.h @@ -6,7 +6,7 @@ #include -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h" diff --git a/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h b/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h index 4eb6823907bcc..4e93f984a88d4 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h @@ -7,7 +7,7 @@ #include #include -#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" +#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "DataFormats/Math/interface/approx_atan2.h" #include "HeterogeneousCore/CUDAUtilities/interface/VecArray.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" 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..51de45237d639 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" From e682d258fc3da6810b255207e6faa9d7b9db616e Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 22:57:46 +0100 Subject: [PATCH 05/40] Allow "if constexpr" in CUDA code --- .../interface/TrackingRecHit2DHeterogeneous.h | 6 +----- .../plugins/CAHitNtupletGeneratorKernelsAlloc.h | 10 +++------- 2 files changed, 4 insertions(+), 12 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h index 73a6daaa4e387..f10495abd2ab8 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h @@ -127,11 +127,7 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous(uint32_t nH m_hitsLayerStart = view->m_hitsLayerStart = reinterpret_cast(get32(n32)); // transfer view - if -#ifndef __CUDACC__ - constexpr -#endif - (std::is_same::value) { + if constexpr (std::is_same::value) { cms::cuda::copyAsync(m_view, view, stream); } else { m_view.reset(view.release()); // NOLINT: std::move() breaks CUDA version diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsAlloc.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsAlloc.h index fb750267f5c37..1c34275d6bbe2 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsAlloc.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsAlloc.h @@ -1,7 +1,7 @@ -#include "CAHitNtupletGeneratorKernels.h" - #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" +#include "CAHitNtupletGeneratorKernels.h" + template <> #ifdef __CUDACC__ void CAHitNtupletGeneratorKernelsGPU::allocateOnGPU(cudaStream_t stream) { @@ -25,11 +25,7 @@ void CAHitNtupletGeneratorKernelsCPU::allocateOnGPU(cudaStream_t stream) { device_hitToTuple_apc_ = (cms::cuda::AtomicPairCounter*)device_storage_.get() + 1; device_nCells_ = (uint32_t*)(device_storage_.get() + 2); - if -#ifndef __CUDACC__ - constexpr -#endif - (std::is_same::value) { + if constexpr (std::is_same::value) { cudaCheck(cudaMemsetAsync(device_nCells_, 0, sizeof(uint32_t), stream)); } else { *device_nCells_ = 0; From ceb57cd02fd236f71eb552c812591850003a6ec9 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 23:09:12 +0100 Subject: [PATCH 06/40] Update comment about m_averageGeometry ownership --- .../TrackingRecHit/interface/TrackingRecHit2DSOAView.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h index 6335a4f9346bf..53297a78a428f 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h @@ -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 From da590430128847af032ba18427be31aa2fd6ec6a Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 23:31:12 +0100 Subject: [PATCH 07/40] Recover missing changes from #29805 --- CalibTracker/SiPixelESProducers/plugins/BuildFile.xml | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml b/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml index 1f063df32a766..4bef676217b72 100644 --- a/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml +++ b/CalibTracker/SiPixelESProducers/plugins/BuildFile.xml @@ -1,13 +1,15 @@ - - + + + + - + From fca227865fe2149131c00d209efc2a114f0404ca Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 14 Dec 2020 23:53:20 +0100 Subject: [PATCH 08/40] Cleanup SiPixelGainCalibrationForHLTGPU Code clean up: - simplify the constructor; - replace commented out messages with LogDebug; - delete commented out code. --- .../src/SiPixelGainCalibrationForHLTGPU.cc | 29 +++++++++---------- 1 file changed, 13 insertions(+), 16 deletions(-) diff --git a/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc b/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc index e4f278c28ec69..980c701fd600f 100644 --- a/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc +++ b/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc @@ -3,6 +3,7 @@ #include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h" #include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h" #include "CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h" +#include "FWCore/MessageLogger/interface/MessageLogger.h" #include "Geometry/CommonDetUnit/interface/GeomDetType.h" #include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" @@ -14,17 +15,16 @@ SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGa auto const& dus = geom.detUnits(); unsigned m_detectors = dus.size(); for (unsigned int i = 1; i < 7; ++i) { - if (geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]) != dus.size() && - dus[geom.offsetDU(GeomDetEnumerators::tkDetEnum[i])]->type().isTrackerStrip()) { - if (geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]) < m_detectors) - m_detectors = geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]); + const auto offset = geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]); + if (offset != dus.size() && dus[offset]->type().isTrackerStrip()) { + if (m_detectors > offset) + m_detectors = offset; } } - /* - std::cout << "caching calibs for " << m_detectors << " pixel detectors of size " << gains.data().size() << std::endl; - std::cout << "sizes " << sizeof(char) << ' ' << sizeof(uint8_t) << ' ' << sizeof(SiPixelGainForHLTonGPU::DecodingStructure) << std::endl; - */ + LogDebug("SiPixelGainCalibrationForHLTGPU") + << "caching calibs for " << m_detectors << " pixel detectors of size " << gains.data().size() << '\n' + << "sizes " << sizeof(char) << ' ' << sizeof(uint8_t) << ' ' << sizeof(SiPixelGainForHLTonGPU::DecodingStructure); cudaCheck(cudaMallocHost((void**)&gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU))); gainForHLTonHost_->v_pedestals = @@ -51,15 +51,12 @@ SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGa gainForHLTonHost_->pedPrecision = static_cast(maxPed - minPed) / nBinsToUseForEncoding; gainForHLTonHost_->gainPrecision = static_cast(maxGain - minGain) / nBinsToUseForEncoding; - /* - std::cout << "precisions g " << gainForHLTonHost_->pedPrecision << ' ' << gainForHLTonHost_->gainPrecision << std::endl; - */ + LogDebug("SiPixelGainCalibrationForHLTGPU") + << "precisions g " << gainForHLTonHost_->pedPrecision << ' ' << gainForHLTonHost_->gainPrecision; // fill the index map auto const& ind = gains.getIndexes(); - /* - std::cout << ind.size() << " " << m_detectors << std::endl; - */ + LogDebug("SiPixelGainCalibrationForHLTGPU") << ind.size() << " " << m_detectors; for (auto i = 0U; i < m_detectors; ++i) { auto p = std::lower_bound( @@ -72,8 +69,8 @@ SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGa assert(p->ibegin != p->iend); assert(p->ncols > 0); gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(p->ibegin, p->iend), p->ncols); - // if (ind[i].detid!=dus[i]->geographicalId()) std::cout << ind[i].detid<<"!="<geographicalId() << std::endl; - // gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(ind[i].ibegin,ind[i].iend), ind[i].ncols); + if (ind[i].detid != dus[i]->geographicalId()) + LogDebug("SiPixelGainCalibrationForHLTGPU") << ind[i].detid << "!=" << dus[i]->geographicalId(); } } From 089d20cad00d2dc0aade618b1d34353a71450b35 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 00:04:35 +0100 Subject: [PATCH 09/40] Rename variables according to the coding rules --- .../src/SiPixelGainCalibrationForHLTGPU.cc | 24 +++++++++---------- .../interface/SiPixelGainForHLTonGPU.h | 23 +++++++++--------- 2 files changed, 23 insertions(+), 24 deletions(-) diff --git a/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc b/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc index 980c701fd600f..66b8d9594353b 100644 --- a/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc +++ b/CalibTracker/SiPixelESProducers/src/SiPixelGainCalibrationForHLTGPU.cc @@ -13,21 +13,21 @@ SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGa : gains_(&gains) { // bizzarre logic (looking for fist strip-det) don't ask auto const& dus = geom.detUnits(); - unsigned m_detectors = dus.size(); + unsigned int n_detectors = dus.size(); for (unsigned int i = 1; i < 7; ++i) { const auto offset = geom.offsetDU(GeomDetEnumerators::tkDetEnum[i]); if (offset != dus.size() && dus[offset]->type().isTrackerStrip()) { - if (m_detectors > offset) - m_detectors = offset; + if (n_detectors > offset) + n_detectors = offset; } } LogDebug("SiPixelGainCalibrationForHLTGPU") - << "caching calibs for " << m_detectors << " pixel detectors of size " << gains.data().size() << '\n' + << "caching calibs for " << n_detectors << " pixel detectors of size " << gains.data().size() << '\n' << "sizes " << sizeof(char) << ' ' << sizeof(uint8_t) << ' ' << sizeof(SiPixelGainForHLTonGPU::DecodingStructure); cudaCheck(cudaMallocHost((void**)&gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU))); - gainForHLTonHost_->v_pedestals = + gainForHLTonHost_->v_pedestals_ = (SiPixelGainForHLTonGPU_DecodingStructure*)this->gains_->data().data(); // so it can be used on CPU as well... // do not read back from the (possibly write-combined) memory buffer @@ -48,17 +48,17 @@ SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGa gainForHLTonHost_->deadFlag_ = 255; gainForHLTonHost_->noisyFlag_ = 254; - gainForHLTonHost_->pedPrecision = static_cast(maxPed - minPed) / nBinsToUseForEncoding; - gainForHLTonHost_->gainPrecision = static_cast(maxGain - minGain) / nBinsToUseForEncoding; + gainForHLTonHost_->pedPrecision_ = static_cast(maxPed - minPed) / nBinsToUseForEncoding; + gainForHLTonHost_->gainPrecision_ = static_cast(maxGain - minGain) / nBinsToUseForEncoding; LogDebug("SiPixelGainCalibrationForHLTGPU") - << "precisions g " << gainForHLTonHost_->pedPrecision << ' ' << gainForHLTonHost_->gainPrecision; + << "precisions g " << gainForHLTonHost_->pedPrecision_ << ' ' << gainForHLTonHost_->gainPrecision_; // fill the index map auto const& ind = gains.getIndexes(); - LogDebug("SiPixelGainCalibrationForHLTGPU") << ind.size() << " " << m_detectors; + LogDebug("SiPixelGainCalibrationForHLTGPU") << ind.size() << " " << n_detectors; - for (auto i = 0U; i < m_detectors; ++i) { + for (auto i = 0U; i < n_detectors; ++i) { auto p = std::lower_bound( ind.begin(), ind.end(), dus[i]->geographicalId().rawId(), SiPixelGainCalibrationForHLT::StrictWeakOrdering()); assert(p != ind.end() && p->detid == dus[i]->geographicalId()); @@ -68,7 +68,7 @@ SiPixelGainCalibrationForHLTGPU::SiPixelGainCalibrationForHLTGPU(const SiPixelGa assert(0 == p->iend % 2); assert(p->ibegin != p->iend); assert(p->ncols > 0); - gainForHLTonHost_->rangeAndCols[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(p->ibegin, p->iend), p->ncols); + gainForHLTonHost_->rangeAndCols_[i] = std::make_pair(SiPixelGainForHLTonGPU::Range(p->ibegin, p->iend), p->ncols); if (ind[i].detid != dus[i]->geographicalId()) LogDebug("SiPixelGainCalibrationForHLTGPU") << ind[i].detid << "!=" << dus[i]->geographicalId(); } @@ -91,7 +91,7 @@ const SiPixelGainForHLTonGPU* SiPixelGainCalibrationForHLTGPU::getGPUProductAsyn cudaCheck(cudaMemcpyAsync( data.gainForHLTonGPU, this->gainForHLTonHost_, sizeof(SiPixelGainForHLTonGPU), cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync(&(data.gainForHLTonGPU->v_pedestals), + cudaCheck(cudaMemcpyAsync(&(data.gainForHLTonGPU->v_pedestals_), &(data.gainDataOnGPU), sizeof(SiPixelGainForHLTonGPU_DecodingStructure*), cudaMemcpyDefault, diff --git a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h index 6326b594e2771..fc228d0207ecf 100644 --- a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h +++ b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h @@ -1,5 +1,5 @@ -#ifndef CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h -#define CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h +#ifndef CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h +#define CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h #include #include @@ -32,8 +32,8 @@ class SiPixelGainForHLTonGPU { inline __host__ __device__ std::pair getPedAndGain( uint32_t moduleInd, int col, int row, bool& isDeadColumn, bool& isNoisyColumn) const { - auto range = rangeAndCols[moduleInd].first; - auto nCols = rangeAndCols[moduleInd].second; + 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; @@ -46,7 +46,7 @@ class SiPixelGainForHLTonGPU { assert(offset < 3088384); assert(0 == offset % 2); - DecodingStructure const* __restrict__ lp = v_pedestals; + DecodingStructure const* __restrict__ lp = v_pedestals_; auto s = lp[offset / 2]; isDeadColumn = (s.ped & 0xFF) == deadFlag_; @@ -55,15 +55,14 @@ class SiPixelGainForHLTonGPU { return std::make_pair(decodePed(s.ped & 0xFF), decodeGain(s.gain & 0xFF)); } - constexpr float decodeGain(unsigned int gain) const { return gain * gainPrecision + minGain_; } - constexpr float decodePed(unsigned int ped) const { return ped * pedPrecision + minPed_; } + constexpr float decodeGain(unsigned int gain) const { return gain * gainPrecision_ + minGain_; } + constexpr float decodePed(unsigned int ped) const { return ped * pedPrecision_ + minPed_; } - DecodingStructure* v_pedestals; - std::pair rangeAndCols[2000]; + DecodingStructure* v_pedestals_; + std::pair rangeAndCols_[2000]; float minPed_, maxPed_, minGain_, maxGain_; - - float pedPrecision, gainPrecision; + float pedPrecision_, gainPrecision_; unsigned int numberOfRowsAveragedOver_; // this is 80!!!! unsigned int nBinsToUseForEncoding_; @@ -71,4 +70,4 @@ class SiPixelGainForHLTonGPU { unsigned int noisyFlag_; }; -#endif // CondFormats_SiPixelObjects_SiPixelGainForHLTonGPU_h +#endif // CondFormats_SiPixelObjects_interface_SiPixelGainForHLTonGPU_h From e307ae2db4e4a91610e0d543d72c2e27ac5ee727 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 00:43:06 +0100 Subject: [PATCH 10/40] Move to LogDebug and other clean up --- .../SiPixelRecHits/interface/PixelCPEBase.h | 60 ++++++++----------- 1 file changed, 24 insertions(+), 36 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h index 4b569438aa130..05e59585ba6ba 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h @@ -1,5 +1,5 @@ -#ifndef RecoLocalTracker_SiPixelRecHits_PixelCPEBase_H -#define RecoLocalTracker_SiPixelRecHits_PixelCPEBase_H 1 +#ifndef RecoLocalTracker_SiPixelRecHits_interface_PixelCPEBase_h +#define RecoLocalTracker_SiPixelRecHits_interface_PixelCPEBase_h 1 //----------------------------------------------------------------------------- // \class PixelCPEBase @@ -11,43 +11,32 @@ // Change to use Generic error & Template calibration from DB - D.Fehling 11/08 //----------------------------------------------------------------------------- +#ifdef EDM_ML_DEBUG +#include +#endif +#include #include #include -#include "TMath.h" -#include "RecoLocalTracker/ClusterParameterEstimator/interface/PixelClusterParameterEstimator.h" -#include "DataFormats/TrackerRecHit2D/interface/SiPixelRecHitQuality.h" +#include +#include "CondFormats/SiPixelObjects/interface/SiPixelGenErrorDBObject.h" +#include "CondFormats/SiPixelObjects/interface/SiPixelLorentzAngle.h" +#include "CondFormats/SiPixelObjects/interface/SiPixelTemplateDBObject.h" +#include "DataFormats/GeometryCommonDetAlgo/interface/MeasurementError.h" +#include "DataFormats/GeometryCommonDetAlgo/interface/MeasurementPoint.h" +#include "DataFormats/GeometrySurface/interface/GloballyPositioned.h" #include "DataFormats/TrackerCommon/interface/TrackerTopology.h" -#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h" +#include "DataFormats/TrackerRecHit2D/interface/SiPixelRecHitQuality.h" +#include "FWCore/MessageLogger/interface/MessageLogger.h" +#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" +#include "FWCore/ParameterSet/interface/ParameterSet.h" #include "Geometry/CommonDetUnit/interface/GeomDetType.h" #include "Geometry/CommonDetUnit/interface/PixelGeomDetUnit.h" #include "Geometry/CommonTopologies/interface/PixelTopology.h" #include "Geometry/CommonTopologies/interface/Topology.h" - -//--- For the configuration: -#include "FWCore/ParameterSet/interface/ParameterSet.h" -#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" - -#include "DataFormats/GeometryCommonDetAlgo/interface/MeasurementPoint.h" -#include "DataFormats/GeometryCommonDetAlgo/interface/MeasurementError.h" -#include "DataFormats/GeometrySurface/interface/GloballyPositioned.h" - -#include "CondFormats/SiPixelObjects/interface/SiPixelLorentzAngle.h" - -// new errors -#include "CondFormats/SiPixelObjects/interface/SiPixelGenErrorDBObject.h" -// old errors -//#include "CondFormats/SiPixelObjects/interface/SiPixelCPEGenericErrorParm.h" - -#include "CondFormats/SiPixelObjects/interface/SiPixelTemplateDBObject.h" - -#include - -#include -#ifdef EDM_ML_DEBUG -#include -#endif +#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h" +#include "RecoLocalTracker/ClusterParameterEstimator/interface/PixelClusterParameterEstimator.h" class RectangularPixelTopology; class MagneticField; @@ -84,7 +73,6 @@ class PixelCPEBase : public PixelClusterParameterEstimator { virtual ~ClusterParam() = default; const SiPixelCluster* theCluster = nullptr; - ; //--- Cluster-level quantities (filled in computeAnglesFrom....) float cotalpha; @@ -145,7 +133,7 @@ class PixelCPEBase : public PixelClusterParameterEstimator { inline ReturnType getParameters(const SiPixelCluster& cl, const GeomDetUnit& det) const override { #ifdef EDM_ML_DEBUG nRecHitsTotal_++; - //std::cout<<" in PixelCPEBase:localParameters(all) - "< Date: Tue, 15 Dec 2020 00:46:59 +0100 Subject: [PATCH 11/40] Rename variables according to the coding rules --- .../plugins/PixelCPEFastESProducer.cc | 39 +++++++++---------- 1 file changed, 18 insertions(+), 21 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc b/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc index 3f7c9aca2a974..facb46360f99c 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc @@ -1,23 +1,20 @@ -#include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h" -#include "MagneticField/Engine/interface/MagneticField.h" -#include "MagneticField/Records/interface/IdealMagneticFieldRecord.h" -#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h" -#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h" -#include "Geometry/Records/interface/TrackerTopologyRcd.h" -#include "DataFormats/TrackerCommon/interface/TrackerTopology.h" -#include "RecoLocalTracker/Records/interface/TkPixelCPERecord.h" -#include "RecoLocalTracker/ClusterParameterEstimator/interface/PixelClusterParameterEstimator.h" +#include +#include -#include "FWCore/Framework/interface/EventSetup.h" +#include "CondFormats/DataRecord/interface/SiPixelGenErrorDBObjectRcd.h" +#include "DataFormats/TrackerCommon/interface/TrackerTopology.h" #include "FWCore/Framework/interface/ESHandle.h" -#include "FWCore/Framework/interface/ModuleFactory.h" #include "FWCore/Framework/interface/ESProducer.h" - -// new record -#include "CondFormats/DataRecord/interface/SiPixelGenErrorDBObjectRcd.h" - -#include -#include +#include "FWCore/Framework/interface/EventSetup.h" +#include "FWCore/Framework/interface/ModuleFactory.h" +#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h" +#include "Geometry/Records/interface/TrackerTopologyRcd.h" +#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h" +#include "MagneticField/Engine/interface/MagneticField.h" +#include "MagneticField/Records/interface/IdealMagneticFieldRecord.h" +#include "RecoLocalTracker/ClusterParameterEstimator/interface/PixelClusterParameterEstimator.h" +#include "RecoLocalTracker/Records/interface/TkPixelCPERecord.h" +#include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h" class PixelCPEFastESProducer : public edm::ESProducer { public: @@ -34,7 +31,7 @@ class PixelCPEFastESProducer : public edm::ESProducer { edm::ESGetToken genErrorDBObjectToken_; edm::ParameterSet pset_; - bool UseErrorsFromTemplates_; + bool useErrorsFromTemplates_; }; using namespace edm; @@ -42,7 +39,7 @@ using namespace edm; PixelCPEFastESProducer::PixelCPEFastESProducer(const edm::ParameterSet& p) : pset_(p) { auto const& myname = p.getParameter("ComponentName"); auto const& magname = p.getParameter("MagneticFieldRecord"); - UseErrorsFromTemplates_ = p.getParameter("UseErrorsFromTemplates"); + useErrorsFromTemplates_ = p.getParameter("UseErrorsFromTemplates"); auto cc = setWhatProduced(this, myname); magfieldToken_ = cc.consumes(magname); @@ -50,7 +47,7 @@ PixelCPEFastESProducer::PixelCPEFastESProducer(const edm::ParameterSet& p) : pse hTTToken_ = cc.consumes(); lorentzAngleToken_ = cc.consumes(edm::ESInputTag("")); lorentzAngleWidthToken_ = cc.consumes(edm::ESInputTag("", "forWidth")); - if (UseErrorsFromTemplates_) { + if (useErrorsFromTemplates_) { genErrorDBObjectToken_ = cc.consumes(); } } @@ -63,7 +60,7 @@ std::unique_ptr PixelCPEFastESProducer::produce( const SiPixelGenErrorDBObject* genErrorDBObjectProduct = nullptr; // Errors take only from new GenError - if (UseErrorsFromTemplates_) { // do only when generrors are needed + if (useErrorsFromTemplates_) { // do only when generrors are needed genErrorDBObjectProduct = &iRecord.get(genErrorDBObjectToken_); //} else { //std::cout<<" pass an empty GenError pointer"< Date: Tue, 15 Dec 2020 00:56:17 +0100 Subject: [PATCH 12/40] Reorder parameters in the autogenerated cfi file --- .../plugins/PixelCPEFastESProducer.cc | 26 +++++++++---------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc b/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc index facb46360f99c..332baabe8842a 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelCPEFastESProducer.cc @@ -75,23 +75,23 @@ std::unique_ptr PixelCPEFastESProducer::produce( } void PixelCPEFastESProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { - // PixelCPEFastESProducer edm::ParameterSetDescription desc; - desc.add("DoLorentz", false); - desc.add("lAWidthFPix", 0); - desc.add("useLAAlignmentOffsets", false); - desc.add("LoadTemplatesFromDB", true); - desc.add("UseErrorsFromTemplates", true); + + // from PixelCPEBase + PixelCPEBase::fillPSetDescription(desc); + + // used by PixelCPEFast desc.add("EdgeClusterErrorX", 50.0); - desc.add("MagneticFieldRecord", edm::ESInputTag()); - desc.add("useLAWidthFromDB", true); - desc.add("TruncatePixelCharge", true); - desc.add("ClusterProbComputationFlag", 0); - desc.add("lAOffset", 0); desc.add("EdgeClusterErrorY", 85.0); + desc.add("UseErrorsFromTemplates", true); + desc.add("TruncatePixelCharge", true); + + // specific to PixelCPEFastESProducer desc.add("ComponentName", "PixelCPEFast"); - desc.add("lAWidthBPix", 0); - desc.add("Alpha2Order", true); + desc.add("MagneticFieldRecord", edm::ESInputTag()); + desc.add("useLAAlignmentOffsets", false); + desc.add("DoLorentz", false); + descriptions.add("PixelCPEFastESProducer", desc); } From 90834a7f03b4b59f8f3cafbabc0541128c62b1b5 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 01:04:26 +0100 Subject: [PATCH 13/40] Rename PixelRecHits.{h,cu} to PixelRecHitGPUKernel.{h,cu} --- .../plugins/{PixelRecHits.cu => PixelRecHitGPUKernel.cu} | 2 +- .../plugins/{PixelRecHits.h => PixelRecHitGPUKernel.h} | 6 +++--- .../SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc | 2 +- 3 files changed, 5 insertions(+), 5 deletions(-) rename RecoLocalTracker/SiPixelRecHits/plugins/{PixelRecHits.cu => PixelRecHitGPUKernel.cu} (98%) rename RecoLocalTracker/SiPixelRecHits/plugins/{PixelRecHits.h => PixelRecHitGPUKernel.h} (84%) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu similarity index 98% rename from RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu rename to RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu index aae20b54e150d..286ba2f4f328d 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu @@ -11,7 +11,7 @@ #include "RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h" -#include "PixelRecHits.h" +#include "PixelRecHitGPUKernel.h" #include "gpuPixelRecHits.h" namespace { diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.h similarity index 84% rename from RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h rename to RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.h index ac02caab3259e..61bc8b58bb7d6 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.h @@ -1,5 +1,5 @@ -#ifndef RecoLocalTracker_SiPixelRecHits_plugins_PixelRecHits_h -#define RecoLocalTracker_SiPixelRecHits_plugins_PixelRecHits_h +#ifndef RecoLocalTracker_SiPixelRecHits_plugins_PixelRecHitGPUKernel_h +#define RecoLocalTracker_SiPixelRecHits_plugins_PixelRecHitGPUKernel_h #include @@ -30,4 +30,4 @@ namespace pixelgpudetails { }; } // namespace pixelgpudetails -#endif // RecoLocalTracker_SiPixelRecHits_plugins_PixelRecHits_h +#endif // RecoLocalTracker_SiPixelRecHits_plugins_PixelRecHitGPUKernel_h diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc index 1312126be32c1..ed651db3f71bf 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc @@ -21,7 +21,7 @@ #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h" #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h" -#include "PixelRecHits.h" // TODO : spit product from kernel +#include "PixelRecHitGPUKernel.h" class SiPixelRecHitCUDA : public edm::global::EDProducer<> { public: From c099c102c01eb0deea6cb9b32aeff63492bd7da5 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 01:28:01 +0100 Subject: [PATCH 14/40] Use gpuClustering::invalidModuleId instead of 9999 --- .../SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 52c02e2510aa3..b3c77f2e17788 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -25,7 +25,7 @@ #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h" #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h" -#include "RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h" +#include "gpuPixelRecHits.h" class SiPixelRecHitSoAFromLegacy : public edm::global::EDProducer<> { public: @@ -215,7 +215,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv if (h - fc < maxHitsInModule) assert(gind == output->view()->detectorIndex(h)); else - assert(9999 == output->view()->detectorIndex(h)); + assert(gpuClustering::invalidModuleId == output->view()->detectorIndex(h)); if (convert2Legacy_) { SiPixelRecHitCollectionNew::FastFiller recHitsOnDetUnit(*legacyOutput, detid); for (auto h = fc; h < lc; ++h) { From c362200150967694667ee48161952144076f833f Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 11:52:19 +0100 Subject: [PATCH 15/40] Update comments about GPU_SMALL_EVENTS --- .../interface/gpuClusteringConstants.h | 8 +++- .../Track/interface/PixelTrackHeterogeneous.h | 2 + .../PixelTriplets/plugins/CAConstants.h | 38 +++++++++---------- 3 files changed, 25 insertions(+), 23 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index 5928d45af7dc8..e9dfed7bca7a6 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -6,18 +6,22 @@ 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; } diff --git a/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h b/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h index e79a32c21daa0..41936b5fc7077 100644 --- 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/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h index 9eea4f528fcdb..d9c3ff70e35ed 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h @@ -10,34 +10,30 @@ #include "HeterogeneousCore/CUDAUtilities/interface/VecArray.h" #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h" -// #define ONLY_PHICUT +//#define ONLY_PHICUT namespace CAConstants { // constants -#ifndef ONLY_PHICUT +#ifdef ONLY_PHICUT + constexpr uint32_t maxNumberOfTuples() { return 48 * 1024; } + constexpr uint32_t maxNumberOfDoublets() { return 2 * 1024 * 1024; } + constexpr uint32_t maxCellsPerHit() { return 8 * 128; } +#else #ifdef GPU_SMALL_EVENTS + // kept for testing and debugging constexpr uint32_t maxNumberOfTuples() { return 3 * 1024; } + constexpr uint32_t maxNumberOfDoublets() { return 128 * 1024; } + constexpr uint32_t maxCellsPerHit() { return 128 / 2; } #else + // tested on MC events with 55-75 pileup events constexpr uint32_t maxNumberOfTuples() { return 24 * 1024; } -#endif -#else - constexpr uint32_t maxNumberOfTuples() { return 48 * 1024; } -#endif - constexpr uint32_t maxNumberOfQuadruplets() { return maxNumberOfTuples(); } -#ifndef ONLY_PHICUT -#ifndef GPU_SMALL_EVENTS constexpr uint32_t maxNumberOfDoublets() { return 512 * 1024; } constexpr uint32_t maxCellsPerHit() { return 128; } -#else - constexpr uint32_t maxNumberOfDoublets() { return 128 * 1024; } - constexpr uint32_t maxCellsPerHit() { return 128 / 2; } -#endif -#else - constexpr uint32_t maxNumberOfDoublets() { return 2 * 1024 * 1024; } - constexpr uint32_t maxCellsPerHit() { return 8 * 128; } #endif +#endif // ONLY_PHICUT constexpr uint32_t maxNumOfActiveDoublets() { return maxNumberOfDoublets() / 8; } + constexpr uint32_t maxNumberOfQuadruplets() { return maxNumberOfTuples(); } constexpr uint32_t maxNumberOfLayerPairs() { return 20; } constexpr uint32_t maxNumberOfLayers() { return 10; } @@ -45,14 +41,14 @@ namespace CAConstants { // types using hindex_type = uint32_t; // FIXME from siPixelRecHitsHeterogeneousProduct - using tindex_type = uint16_t; // for tuples + using tindex_type = uint16_t; // for tuples -#ifndef ONLY_PHICUT - using CellNeighbors = cms::cuda::VecArray; - using CellTracks = cms::cuda::VecArray; -#else +#ifdef ONLY_PHICUT using CellNeighbors = cms::cuda::VecArray; using CellTracks = cms::cuda::VecArray; +#else + using CellNeighbors = cms::cuda::VecArray; + using CellTracks = cms::cuda::VecArray; #endif using CellNeighborsVector = cms::cuda::SimpleVector; From f6924bb184f1a9b34cddd602d1cd84891a9ea98b Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 15:09:04 +0100 Subject: [PATCH 16/40] Consistently use gpuClustering::maxNumModules --- .../src/TrackingRecHit2DHeterogeneous.cc | 4 ++-- .../interface/SiPixelGainForHLTonGPU.h | 3 ++- .../plugins/SiPixelDigisClustersFromSoA.cc | 3 ++- .../SiPixelClusterizer/test/gpuClustering_t.h | 3 ++- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 12 ++++++------ .../PixelTriplets/plugins/gpuPixelDoubletsAlgos.h | 4 ++-- .../TrackerHitAssociation/plugins/ClusterSLOnGPU.cu | 5 +++-- 7 files changed, 19 insertions(+), 15 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc index d4bf1b500e216..7df49b1c9f780 100644 --- a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc +++ b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc @@ -13,7 +13,7 @@ cms::cuda::host::unique_ptr TrackingRecHit2DCUDA::localCoordToHostAsync template <> cms::cuda::host::unique_ptr TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cudaStream_t stream) const { - auto ret = cms::cuda::make_host_unique(2001, stream); - cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, 4 * 2001, cudaMemcpyDefault, stream)); + auto ret = cms::cuda::make_host_unique(gpuClustering::maxNumModules + 1, stream); + cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClustering::maxNumModules + 1), cudaMemcpyDefault, stream)); return ret; } diff --git a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h index fc228d0207ecf..aa5a127927b90 100644 --- a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h +++ b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h @@ -16,6 +16,7 @@ #endif // __device__ #endif // __CUDACC__ +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" struct SiPixelGainForHLTonGPU_DecodingStructure { @@ -59,7 +60,7 @@ class SiPixelGainForHLTonGPU { constexpr float decodePed(unsigned int ped) const { return ped * pedPrecision_ + minPed_; } DecodingStructure* v_pedestals_; - std::pair rangeAndCols_[2000]; + std::pair rangeAndCols_[gpuClustering::maxNumModules]; float minPed_, maxPed_, minGain_, maxGain_; float pedPrecision_, gainPrecision_; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index dbbc5c4b03284..0685a1d1abed7 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -1,3 +1,4 @@ +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "DataFormats/Common/interface/DetSetVector.h" #include "DataFormats/Common/interface/Handle.h" #include "DataFormats/DetId/interface/DetId.h" @@ -84,7 +85,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con auto collection = std::make_unique>(); auto outputClusters = std::make_unique(); - outputClusters->reserve(2000, nDigis / 4); + outputClusters->reserve(gpuClustering::maxNumModules, nDigis / 4); edm::DetSet* detDigis = nullptr; for (uint32_t i = 0; i < nDigis; i++) { diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h index e3e5f17604df0..02611ab1cac1d 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h +++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h @@ -26,7 +26,8 @@ int main(void) { using namespace gpuClustering; - int numElements = 256 * 2000; + constexpr int numElements = 256 * maxNumModules; + // these in reality are already on GPU auto h_id = std::make_unique(numElements); auto h_x = std::make_unique(numElements); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index b3c77f2e17788..4f4e1e4113564 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -113,8 +113,8 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv HitModuleStart moduleStart_; // index of the first pixel of each module HitModuleStart clusInModule_; memset(&clusInModule_, 0, sizeof(HitModuleStart)); // needed?? - assert(2001 == clusInModule_.size()); - assert(0 == clusInModule_[2000]); + assert(gpuClustering::maxNumModules + 1 == clusInModule_.size()); + assert(0 == clusInModule_[gpuClustering::maxNumModules]); uint32_t moduleId_; moduleStart_[1] = 0; // we run sequentially.... @@ -128,7 +128,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv DetId detIdObject(detid); const GeomDetUnit* genericDet = geom_->idToDetUnit(detIdObject); auto gind = genericDet->index(); - assert(gind < 2000); + assert(gind < gpuClustering::maxNumModules); auto const nclus = DSViter->size(); clusInModule_[gind] = nclus; numberOfClusters += nclus; @@ -136,7 +136,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv hitsModuleStart[0] = 0; for (int i = 1, n = clusInModule_.size(); i < n; ++i) hitsModuleStart[i] = hitsModuleStart[i - 1] + clusInModule_[i - 1]; - assert(numberOfClusters == int(hitsModuleStart[2000])); + assert(numberOfClusters == int(hitsModuleStart[gpuClustering::maxNumModules])); // output SoA auto output = std::make_unique(numberOfClusters, &cpeView, hitsModuleStart, nullptr); @@ -149,7 +149,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv } if (convert2Legacy_) - legacyOutput->reserve(2000, numberOfClusters); + legacyOutput->reserve(gpuClustering::maxNumModules, numberOfClusters); int numberOfDetUnits = 0; int numberOfHits = 0; @@ -159,7 +159,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv DetId detIdObject(detid); const GeomDetUnit* genericDet = geom_->idToDetUnit(detIdObject); auto const gind = genericDet->index(); - assert(gind < 2000); + assert(gind < gpuClustering::maxNumModules); const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); auto const nclus = DSViter->size(); diff --git a/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h b/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h index 4e93f984a88d4..d055c8b7cb867 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h @@ -105,7 +105,7 @@ namespace gpuPixelDoublets { // found hit corresponding to our cuda thread, now do the job auto mi = hh.detectorIndex(i); - if (mi > 2000) + if (mi > gpuClustering::maxNumModules) continue; // invalid /* maybe clever, not effective when zoCut is on @@ -201,7 +201,7 @@ namespace gpuPixelDoublets { assert(oi >= offsets[outer]); assert(oi < offsets[outer + 1]); auto mo = hh.detectorIndex(oi); - if (mo > 2000) + if (mo > gpuClustering::maxNumModules) continue; // invalid if (doZ0Cut && z0cutoff(oi)) diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu index c06d6d254bad3..0aab26d9cc091 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu +++ b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu @@ -22,6 +22,7 @@ __global__ void simLink(const SiPixelDigisCUDA::DeviceConstView* dd, uint32_t n) { 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; @@ -32,12 +33,12 @@ __global__ void simLink(const SiPixelDigisCUDA::DeviceConstView* dd, auto id = dd->moduleInd(i); 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}}; From f217a050d121cef3b06aa07f9d4d510e2660c268 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 16:24:58 +0100 Subject: [PATCH 17/40] Extend the SiPixelCluster constructor --- .../SiPixelCluster/interface/SiPixelCluster.h | 13 ++++++++----- .../plugins/SiPixelDigisClustersFromSoA.cc | 3 +-- 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/DataFormats/SiPixelCluster/interface/SiPixelCluster.h b/DataFormats/SiPixelCluster/interface/SiPixelCluster.h index 5dfb8671c0a38..453d41555a65d 100644 --- a/DataFormats/SiPixelCluster/interface/SiPixelCluster.h +++ b/DataFormats/SiPixelCluster/interface/SiPixelCluster.h @@ -69,19 +69,22 @@ class SiPixelCluster { static constexpr unsigned int MAXSPAN = 255; static constexpr unsigned int MAXPOS = 2047; + static constexpr uint16_t invalidClusterId = std::numeric_limits::max(); + /** Construct from a range of digis that form a cluster and from * a DetID. The range is assumed to be non-empty. */ - SiPixelCluster() {} + SiPixelCluster() = default; SiPixelCluster(unsigned int isize, uint16_t const* adcs, uint16_t const* xpos, uint16_t const* ypos, - uint16_t const xmin, - uint16_t const ymin) - : thePixelOffset(2 * isize), thePixelADC(adcs, adcs + isize) { + uint16_t xmin, + uint16_t ymin, + uint16_t id = invalidClusterId) + : thePixelOffset(2 * isize), thePixelADC(adcs, adcs + isize), theOriginalClusterId(id) { uint16_t maxCol = 0; uint16_t maxRow = 0; for (unsigned int i = 0; i != isize; ++i) { @@ -203,7 +206,7 @@ class SiPixelCluster { uint8_t thePixelRowSpan = 0; // Span pixel index in the x direction (low edge). uint8_t thePixelColSpan = 0; // Span pixel index in the y direction (left edge). - uint16_t theOriginalClusterId = std::numeric_limits::max(); + uint16_t theOriginalClusterId = invalidClusterId; float err_x = -99999.9f; float err_y = -99999.9f; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index 0685a1d1abed7..af260a46cb8bc 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -114,8 +114,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con edm::LogWarning("SiPixelDigisClustersFromSoA") << "cluster below charge Threshold " << "Layer/DetId/clusId " << layer << '/' << detId << '/' << ic << " size/charge " << acluster.isize << '/' << acluster.charge; - SiPixelCluster cluster(acluster.isize, acluster.adc, acluster.x, acluster.y, acluster.xmin, acluster.ymin); - cluster.setOriginalId(ic); + SiPixelCluster cluster(acluster.isize, acluster.adc, acluster.x, acluster.y, acluster.xmin, acluster.ymin, ic); ++totCluseFilled; // std::cout << "putting in this cluster " << ic << " " << cluster.charge() << " " << cluster.pixelADC().size() << endl; // sort by row (x) From 1894353108846a13e19a9dc6f599f04cfef1a221 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 22:58:17 +0100 Subject: [PATCH 18/40] Remove forwarding header file --- .../plugins/SiPixelClusterProducer.cc | 17 ++++++++--------- .../SiPixelClusterizer/plugins/gpuCalibPixel.h | 2 -- .../plugins/gpuClusterChargeCut.h | 3 +-- .../SiPixelClusterizer/plugins/gpuClustering.h | 3 +-- .../plugins/gpuClusteringConstants.h | 6 ------ .../plugins/PixelRecHitGPUKernel.cu | 2 +- 6 files changed, 11 insertions(+), 22 deletions(-) delete mode 100644 RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc index 02678c999a036..5a8ea9e772163 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc @@ -124,10 +124,10 @@ void SiPixelClusterProducer::produce(edm::Event& e, const edm::EventSetup& es) { output->shrink_to_fit(); // set sequential identifier (this is a const interface, but we need to set it after the sorting) - for (auto DSViter = output->begin(); DSViter != output->end(); DSViter++) { + for (auto& clusters : *output) { uint16_t id = 0; - for (auto& clust : *DSViter) { - const_cast(clust).setOriginalId(id++); + for (auto& cluster : clusters) { + cluster.setOriginalId(id++); } } e.put(tPutPixelClusters, std::move(output)); @@ -161,15 +161,14 @@ void SiPixelClusterProducer::run(const T& input, int numberOfClusters = 0; // Iterate on detector units - typename T::const_iterator DSViter = input.begin(); - for (; DSViter != input.end(); DSViter++) { + for (auto const& dsv : input) { ++numberOfDetUnits; // LogDebug takes very long time, get rid off. - //LogDebug("SiStripClusterizer") << "[SiPixelClusterProducer::run] DetID" << DSViter->id; + //LogDebug("SiStripClusterizer") << "[SiPixelClusterProducer::run] DetID" << dsv.id; std::vector badChannels; - DetId detIdObject(DSViter->detId()); + DetId detIdObject(dsv.detId()); // Comment: At the moment the clusterizer depends on geometry // to access information as the pixel topology (number of columns @@ -185,8 +184,8 @@ void SiPixelClusterProducer::run(const T& input, { // Produce clusters for this DetUnit and store them in // a DetSet - edmNew::DetSetVector::FastFiller spc(output, DSViter->detId()); - clusterizer_->clusterizeDetUnit(*DSViter, pixDet, tTopo_, badChannels, spc); + edmNew::DetSetVector::FastFiller spc(output, dsv.detId()); + clusterizer_->clusterizeDetUnit(dsv, pixDet, tTopo_, badChannels, spc); if (spc.empty()) { spc.abort(); } else { diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h index 572a482144667..c21c792f39c30 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h @@ -8,8 +8,6 @@ #include "CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" -#include "gpuClusteringConstants.h" - namespace gpuCalibPixel { using gpuClustering::invalidModuleId; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h index 8a90134488ec5..d9520da80b695 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -4,11 +4,10 @@ #include #include +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "HeterogeneousCore/CUDAUtilities/interface/prefixScan.h" -#include "gpuClusteringConstants.h" - namespace gpuClustering { __global__ void clusterChargeCut( diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h index 508138f5fbb57..56bbab083483d 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h @@ -4,12 +4,11 @@ #include #include +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h" #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" -#include "gpuClusteringConstants.h" - namespace gpuClustering { #ifdef GPU_DEBUG diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h deleted file mode 100644 index 0d5803a760bf6..0000000000000 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h +++ /dev/null @@ -1,6 +0,0 @@ -#ifndef RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusteringConstants_h -#define RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusteringConstants_h - -#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" - -#endif // RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusteringConstants_h diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu index 286ba2f4f328d..f75d5e3b3bef7 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHitGPUKernel.cu @@ -6,10 +6,10 @@ #include // CMSSW headers +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h" -#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusteringConstants.h" #include "PixelRecHitGPUKernel.h" #include "gpuPixelRecHits.h" From 3f06c72ced5d88112de0c1c78e66223540b865c7 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 15 Dec 2020 23:55:59 +0100 Subject: [PATCH 19/40] Added comments --- EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc index 7a49646d7a9a1..f8cb585f604f3 100644 --- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc +++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc @@ -60,6 +60,7 @@ SiPixelDigiErrorsFromSoA::SiPixelDigiErrorsFromSoA(const edm::ParameterSet& iCon void SiPixelDigiErrorsFromSoA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; desc.add("digiErrorSoASrc", edm::InputTag("siPixelDigiErrorsSoA")); + // the configuration parameters here are named following those in SiPixelRawToDigi desc.add("CablingMapLabel", "")->setComment("CablingMap label"); desc.add("UsePhase1", false)->setComment("## Use phase1"); desc.add>("ErrorList", std::vector{29}) From 9dc393168636f3c4e48986f7739c554f75cdd1c1 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 00:46:24 +0100 Subject: [PATCH 20/40] General code cleanup General clean up of the pixel local reconstructon code: - improve comments, remove commented out code and obsolete comments - replace std::cout with LogDebug - update variable names to follow the coding rules --- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 51 +++++----- .../SiPixelRecHits/src/PixelCPEFast.cc | 92 +++++++------------ 2 files changed, 58 insertions(+), 85 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 4f4e1e4113564..26688b1c10003 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -1,6 +1,5 @@ #include -// hack waiting for if constexpr #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h" #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h" @@ -90,21 +89,24 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv iEvent.getByToken(clusterToken_, hclusters); auto const& input = *hclusters; - // yes a unique ptr of a unique ptr so edm is happy and the pointer stay still... + // allocate a buffer for the indices of the clusters auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); + // hitsModuleStart is a non-owning pointer to the buffer auto hitsModuleStart = hmsp.get(); - auto hms = std::make_unique(std::move(hmsp)); // hmsp is gone - iEvent.put(tokenModuleStart_, std::move(hms)); // hms is gone! hitsModuleStart still alive and kicking... + // wrap the buffer in a HostProduct + auto hms = std::make_unique(std::move(hmsp)); + // move the HostProduct to the Event, without reallocating the buffer or affecting hitsModuleStart + iEvent.put(tokenModuleStart_, std::move(hms)); // legacy output auto legacyOutput = std::make_unique(); // storage - std::vector xx_; - std::vector yy_; - std::vector adc_; - std::vector moduleInd_; - std::vector clus_; + std::vector xx; + std::vector yy; + std::vector adc; + std::vector moduleInd; + std::vector clus; std::vector, SiPixelCluster>> clusterRef; @@ -170,19 +172,19 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv auto const fc = hitsModuleStart[gind]; auto const lc = hitsModuleStart[gind + 1]; assert(lc > fc); - // std::cout << "in det " << gind << ": conv " << nclus << " hits from " << DSViter->size() << " legacy clusters" - // <<' '<< fc <<','<size() << " legacy clusters" << ' ' << fc << ',' << lc; assert((lc - fc) == nclus); if (nclus > maxHitsInModule) printf( "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nclus, gind, maxHitsInModule); // fill digis - xx_.clear(); - yy_.clear(); - adc_.clear(); - moduleInd_.clear(); - clus_.clear(); + xx.clear(); + yy.clear(); + adc.clear(); + moduleInd.clear(); + clus.clear(); clusterRef.clear(); moduleId_ = gind; uint32_t ic = 0; @@ -191,11 +193,11 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv assert(clust.size() > 0); for (int i = 0, nd = clust.size(); i < nd; ++i) { auto px = clust.pixel(i); - xx_.push_back(px.x); - yy_.push_back(px.y); - adc_.push_back(px.adc); - moduleInd_.push_back(gind); - clus_.push_back(ic); + xx.push_back(px.x); + yy.push_back(px.y); + adc.push_back(px.adc); + moduleInd.push_back(gind); + clus.push_back(ic); ++ndigi; } assert(clust.originalId() == ic); // make sure hits and clus are in sync @@ -204,10 +206,10 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv ic++; } assert(nclus == ic); - assert(clus_.size() == ndigi); + assert(clus.size() == ndigi); numberOfHits += nclus; // filled creates view - SiPixelDigisCUDA::DeviceConstView digiView{xx_.data(), yy_.data(), adc_.data(), moduleInd_.data(), clus_.data()}; + SiPixelDigisCUDA::DeviceConstView digiView{xx.data(), yy.data(), adc.data(), moduleInd.data(), clus.data()}; assert(digiView.adc(0) != 0); // we run on blockId.x==0 gpuPixelRecHits::getHits(&cpeView, &bsHost, &digiView, ndigi, &clusterView, output->view()); @@ -240,7 +242,8 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv cms::cuda::fillManyFromVector( output->phiBinner(), 10, output->iphi(), output->hitsLayerStart(), numberOfHits, 256, nullptr); - // std::cout << "created HitSoa for " << numberOfClusters << " clusters in " << numberOfDetUnits << " Dets" << std::endl; + LogDebug("SiPixelRecHitSoAFromLegacy") << "created HitSoa for " << numberOfClusters << " clusters in " + << numberOfDetUnits << " Dets"; iEvent.put(std::move(output)); if (convert2Legacy_) iEvent.put(std::move(legacyOutput)); diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc index f3b3f308fa9d3..774727cf91d9c 100644 --- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc +++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc @@ -1,5 +1,3 @@ -#include - #include #include @@ -112,7 +110,8 @@ void PixelCPEFast::fillParamsForGpu() { m_commonParamsGPU.thePitchX = m_DetParams[0].thePitchX; m_commonParamsGPU.thePitchY = m_DetParams[0].thePitchY; - // std::cout << "pitch & thickness " << m_commonParamsGPU.thePitchX << ' ' << m_commonParamsGPU.thePitchY << " " << m_commonParamsGPU.theThicknessB << ' ' << m_commonParamsGPU.theThicknessE << std::endl; + LogDebug("PixelCPEFast") << "pitch & thickness " << m_commonParamsGPU.thePitchX << ' ' << m_commonParamsGPU.thePitchY + << " " << m_commonParamsGPU.theThicknessB << ' ' << m_commonParamsGPU.theThicknessE; // zero average geometry memset(&m_averageGeometry, 0, sizeof(pixelCPEforGPU::AverageGeometry)); @@ -132,30 +131,28 @@ void PixelCPEFast::fillParamsForGpu() { assert(p.theDet->index() == int(i)); assert(m_commonParamsGPU.thePitchY == p.thePitchY); assert(m_commonParamsGPU.thePitchX == p.thePitchX); - //assert(m_commonParamsGPU.theThickness==p.theThickness); g.isBarrel = GeomDetEnumerators::isBarrel(p.thePart); g.isPosZ = p.theDet->surface().position().z() > 0; g.layer = ttopo_.layer(p.theDet->geographicalId()); g.index = i; // better be! g.rawId = p.theDet->geographicalId(); - assert((g.isBarrel ? m_commonParamsGPU.theThicknessB : m_commonParamsGPU.theThicknessE) == p.theThickness); - //if (m_commonParamsGPU.theThickness!=p.theThickness) - // std::cout << i << (g.isBarrel ? "B " : "E ") << m_commonParamsGPU.theThickness<<"!="<geographicalId()); if (oldLayer != g.layer) { oldLayer = g.layer; - // std::cout << "new layer at " << i << (g.isBarrel ? " B " : (g.isPosZ ? " E+ " : " E- ")) << g.layer << " starting at " << g.rawId << std::endl; - // std::cout << "old layer had " << nl << " ladders" << std::endl; + LogDebug("PixelCPEFast") << "new layer at " << i << (g.isBarrel ? " B " : (g.isPosZ ? " E+ " : " E- ")) + << g.layer << " starting at " << g.rawId << '\n' + << "old layer had " << nl << " ladders"; nl = 0; } if (oldLadder != ladder) { oldLadder = ladder; - // std::cout << "new ladder at " << i << (g.isBarrel ? " B " : (g.isPosZ ? " E+ " : " E- ")) << ladder << " starting at " << g.rawId << std::endl; - // std::cout << "old ladder ave z,r,p mz " << zl/8.f << " " << rl/8.f << " " << pl/8.f << ' ' << miz << ' ' << mxz << std::endl; + LogDebug("PixelCPEFast") << "new ladder at " << i << (g.isBarrel ? " B " : (g.isPosZ ? " E+ " : " E- ")) + << ladder << " starting at " << g.rawId << '\n' + << "old ladder ave z,r,p mz " << zl / 8.f << " " << rl / 8.f << " " << pl / 8.f << ' ' + << miz << ' ' << mxz; rl = 0; zl = 0; pl = 0; @@ -200,16 +197,15 @@ void PixelCPEFast::fillParamsForGpu() { if (lape.invalid()) lape = LocalError(); // zero.... -#ifdef DUMP_ERRORS +#ifdef EDM_ML_DEBUG auto m = 10000.f; for (float qclus = 15000; qclus < 35000; qclus += 15000) { errorFromTemplates(p, cp, qclus); - - std::cout << i << ' ' << qclus << ' ' << cp.pixmx << ' ' << m * cp.sigmax << ' ' << m * cp.sx1 << ' ' - << m * cp.sx2 << ' ' << m * cp.sigmay << ' ' << m * cp.sy1 << ' ' << m * cp.sy2 << std::endl; + LogDebug("PixelCPEFast") << i << ' ' << qclus << ' ' << cp.pixmx << ' ' << m * cp.sigmax << ' ' << m * cp.sx1 + << ' ' << m * cp.sx2 << ' ' << m * cp.sigmay << ' ' << m * cp.sy1 << ' ' << m * cp.sy2; } - std::cout << i << ' ' << m * std::sqrt(lape.xx()) << ' ' << m * std::sqrt(lape.yy()) << std::endl; -#endif + LogDebug("PixelCPEFast") << i << ' ' << m * std::sqrt(lape.xx()) << ' ' << m * std::sqrt(lape.yy()); +#endif // EDM_ML_DEBUG errorFromTemplates(p, cp, 20000.f); g.pixmx = std::max(0, cp.pixmx); @@ -221,35 +217,6 @@ void PixelCPEFast::fillParamsForGpu() { g.sy[1] = cp.sy1; g.sy[2] = cp.sy2; - /* - // from run1?? - if (i<96) { - g.sx[0] = 0.00120; - g.sx[1] = 0.00115; - g.sx[2] = 0.0050; - - g.sy[0] = 0.00210; - g.sy[1] = 0.00375; - g.sy[2] = 0.0085; - } else if (g.isBarrel) { - g.sx[0] = 0.00120; - g.sx[1] = 0.00115; - g.sx[2] = 0.0050; - - g.sy[0] = 0.00210; - g.sy[1] = 0.00375; - g.sy[2] = 0.0085; - } else { - g.sx[0] = 0.0020; - g.sx[1] = 0.0020; - g.sx[2] = 0.0050; - - g.sy[0] = 0.0021; - g.sy[1] = 0.0021; - g.sy[2] = 0.0085; - } - */ - for (int i = 0; i < 3; ++i) { g.sx[i] = std::sqrt(g.sx[i] * g.sx[i] + lape.xx()); g.sy[i] = std::sqrt(g.sy[i] * g.sy[i] + lape.yy()); @@ -269,7 +236,7 @@ void PixelCPEFast::fillParamsForGpu() { aveGeom.ladderMaxZ[il] = std::max(aveGeom.ladderMaxZ[il], z); aveGeom.ladderX[il] += 0.125f * g.frame.x(); aveGeom.ladderY[il] += 0.125f * g.frame.y(); - aveGeom.ladderR[il] += 0.125 * sqrt(g.frame.x() * g.frame.x() + g.frame.y() * g.frame.y()); + aveGeom.ladderR[il] += 0.125f * sqrt(g.frame.x() * g.frame.x() + g.frame.y() * g.frame.y()); } assert(il + 1 == int(phase1PixelTopology::numberOfLaddersInBarrel)); // add half_module and tollerance @@ -293,13 +260,16 @@ void PixelCPEFast::fillParamsForGpu() { aveGeom.endCapZ[0] -= 1.5f; aveGeom.endCapZ[1] += 1.5f; - /* - for (int jl=0, nl=phase1PixelTopology::numberOfLaddersInBarrel; jl Date: Wed, 16 Dec 2020 08:08:49 +0100 Subject: [PATCH 21/40] General code cleanup General clean up of the pixel local reconstructon code: - improve comments, remove commented out code and obsolete comments - replace std::cout with LogDebug - update variable names to follow the coding rules - reuse existing constants - remove some obsolete code --- .../src/SiPixelROCsStatusAndMappingWrapper.cc | 16 +++++------ .../interface/SiPixelROCsStatusAndMapping.h | 2 +- .../plugins/SiPixelRawToClusterGPUKernel.cu | 27 +++++++++---------- .../plugins/SiPixelRawToClusterGPUKernel.h | 15 +---------- .../plugins/gpuClustering.h | 1 - .../SiPixelRecHits/src/PixelCPEFast.cc | 4 +-- 6 files changed, 24 insertions(+), 41 deletions(-) diff --git a/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc b/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc index 1470ad6825b86..2437696656d25 100644 --- a/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc +++ b/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc @@ -43,7 +43,7 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe cablingMapHost->link[index] = link; cablingMapHost->roc[index] = roc; if (pixelRoc != nullptr) { - cablingMapHost->RawId[index] = pixelRoc->rawId(); + cablingMapHost->rawId[index] = pixelRoc->rawId(); cablingMapHost->rocInDet[index] = pixelRoc->idInDetUnit(); modToUnpDefault[index] = false; if (badPixelInfo != nullptr) @@ -51,7 +51,7 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe else cablingMapHost->badRocs[index] = false; } else { // store some dummy number - cablingMapHost->RawId[index] = 9999; + cablingMapHost->rawId[index] = 9999; cablingMapHost->rocInDet[index] = 9999; cablingMapHost->badRocs[index] = true; modToUnpDefault[index] = true; @@ -62,7 +62,7 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe } // end of FED loop // Given FedId, Link and idinLnk; use the following formula - // to get the RawId and idinDU + // to get the rawId and idinDU // index = (FedID-1200) * MAX_LINK* MAX_ROC + (Link-1)* MAX_ROC + idinLnk; // where, MAX_LINK = 48, MAX_ROC = 8 for Phase1 as mentioned Danek's email // FedID varies between 1200 to 1338 (In total 108 FED's) @@ -70,15 +70,15 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe // idinLnk varies between 1 to 8 for (int i = 1; i < index; i++) { - if (cablingMapHost->RawId[i] == 9999) { + if (cablingMapHost->rawId[i] == 9999) { cablingMapHost->moduleId[i] = 9999; } else { /* - std::cout << cablingMapHost->RawId[i] << std::endl; + std::cout << cablingMapHost->rawId[i] << std::endl; */ - auto gdet = trackerGeom.idToDetUnit(cablingMapHost->RawId[i]); + auto gdet = trackerGeom.idToDetUnit(cablingMapHost->rawId[i]); if (!gdet) { - LogDebug("SiPixelROCsStatusAndMapping") << " Not found: " << cablingMapHost->RawId[i] << std::endl; + LogDebug("SiPixelROCsStatusAndMapping") << " Not found: " << cablingMapHost->rawId[i] << std::endl; continue; } cablingMapHost->moduleId[i] = gdet->index(); @@ -89,7 +89,7 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe << i << std::setw(20) << cablingMapHost->fed[i] << std::setw(20) << cablingMapHost->link[i] << std::setw(20) << cablingMapHost->roc[i] << std::endl; LogDebug("SiPixelROCsStatusAndMapping") - << i << std::setw(20) << cablingMapHost->RawId[i] << std::setw(20) << cablingMapHost->rocInDet[i] + << i << std::setw(20) << cablingMapHost->rawId[i] << std::setw(20) << cablingMapHost->rocInDet[i] << std::setw(20) << cablingMapHost->moduleId[i] << std::endl; LogDebug("SiPixelROCsStatusAndMapping") << i << std::setw(20) << (bool)cablingMapHost->badRocs[i] << std::setw(20) << std::endl; diff --git a/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h b/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h index df5b8b24b70dc..a0771aaefb366 100644 --- a/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h +++ b/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h @@ -16,7 +16,7 @@ struct SiPixelROCsStatusAndMapping { alignas(128) unsigned int fed[pixelgpudetails::MAX_SIZE]; alignas(128) unsigned int link[pixelgpudetails::MAX_SIZE]; alignas(128) unsigned int roc[pixelgpudetails::MAX_SIZE]; - alignas(128) unsigned int RawId[pixelgpudetails::MAX_SIZE]; + alignas(128) unsigned int rawId[pixelgpudetails::MAX_SIZE]; alignas(128) unsigned int rocInDet[pixelgpudetails::MAX_SIZE]; alignas(128) unsigned int moduleId[pixelgpudetails::MAX_SIZE]; alignas(128) unsigned char badRocs[pixelgpudetails::MAX_SIZE]; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index b1ab47be7dad6..32143e7e5c0a2 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -21,6 +21,8 @@ // CMSSW includes #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h" +#include "DataFormats/FEDRawData/interface/FEDNumbering.h" +#include "DataFormats/TrackerCommon/interface/TrackerTopology.h" #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" @@ -46,7 +48,7 @@ namespace pixelgpudetails { const cms_uint32_t *src, unsigned int length) { std::memcpy(word_.get() + wordCounterGPU, src, sizeof(cms_uint32_t) * length); - std::memset(fedId_.get() + wordCounterGPU / 2, fedId - 1200, length / 2); + std::memset(fedId_.get() + wordCounterGPU / 2, fedId - FEDNumbering::MINSiPixeluTCAFEDID, length / 2); } //////////////////// @@ -59,7 +61,7 @@ namespace pixelgpudetails { __device__ uint32_t getADC(uint32_t ww) { return ((ww >> pixelgpudetails::ADC_shift) & pixelgpudetails::ADC_mask); } - __device__ bool isBarrel(uint32_t rawId) { return (1 == ((rawId >> 25) & 0x7)); } + __device__ bool isBarrel(uint32_t rawId) { return (PixelSubdetector::PixelBarrel == ((rawId >> DetId::kSubdetOffset) & DetId::kSubdetMask)); } __device__ pixelgpudetails::DetIdGPU getRawId(const SiPixelROCsStatusAndMapping *cablingMap, uint8_t fed, @@ -67,7 +69,7 @@ namespace pixelgpudetails { uint32_t roc) { uint32_t index = fed * MAX_LINK * MAX_ROC + (link - 1) * MAX_ROC + roc; pixelgpudetails::DetIdGPU detId = { - cablingMap->RawId[index], cablingMap->rocInDet[index], cablingMap->moduleId[index]}; + cablingMap->rawId[index], cablingMap->rocInDet[index], cablingMap->moduleId[index]}; return detId; } @@ -137,7 +139,7 @@ namespace pixelgpudetails { uint32_t gRow = rowOffset + slopeRow * local.row; uint32_t gCol = colOffset + slopeCol * local.col; - //printf("Inside frameConversion row: %u, column: %u\n", gRow, gCol); + // inside frameConversion row: gRow, column: gCol pixelgpudetails::Pixel global = {gRow, gCol}; return global; } @@ -145,8 +147,6 @@ namespace pixelgpudetails { __device__ uint8_t conversionError(uint8_t fedId, uint8_t status, bool debug = false) { uint8_t errorType = 0; - // debug = true; - switch (status) { case (1): { if (debug) @@ -181,11 +181,8 @@ namespace pixelgpudetails { } __device__ bool rocRowColIsValid(uint32_t rocRow, uint32_t rocCol) { - uint32_t numRowsInRoc = 80; - uint32_t numColsInRoc = 52; - - /// row and collumn in ROC representation - return ((rocRow < numRowsInRoc) & (rocCol < numColsInRoc)); + /// row and column in ROC representation + return ((rocRow < pixelgpudetails::numRowsInRoc) & (rocCol < pixelgpudetails::numColsInRoc)); } __device__ bool dcolIsValid(uint32_t dcol, uint32_t pxid) { return ((dcol < 26) & (2 <= pxid) & (pxid < 162)); } @@ -287,7 +284,7 @@ namespace pixelgpudetails { //cabling.pxid = 2; uint32_t roc = 1; uint32_t link = (errWord >> pixelgpudetails::LINK_shift) & pixelgpudetails::LINK_mask; - uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).RawId; + uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; if (rID_temp != 9999) rID = rID_temp; break; @@ -323,7 +320,7 @@ namespace pixelgpudetails { //cabling.pxid = 2; uint32_t roc = 1; uint32_t link = chanNmbr; - uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).RawId; + uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; if (rID_temp != 9999) rID = rID_temp; break; @@ -334,7 +331,7 @@ namespace pixelgpudetails { //cabling.pxid = 2; uint32_t roc = (errWord >> pixelgpudetails::ROC_shift) & pixelgpudetails::ROC_mask; uint32_t link = (errWord >> pixelgpudetails::LINK_shift) & pixelgpudetails::LINK_mask; - uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).RawId; + uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; if (rID_temp != 9999) rID = rID_temp; break; @@ -397,7 +394,7 @@ namespace pixelgpudetails { continue; } - uint32_t rawId = detId.RawId; + uint32_t rawId = detId.rawId; uint32_t rocIdInDetUnit = detId.rocInDet; bool barrel = isBarrel(rawId); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index e06ba8ce735aa..75eeab2606dd5 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h @@ -75,7 +75,7 @@ namespace pixelgpudetails { const uint32_t OMIT_ERR_mask = ~(~uint32_t(0) << OMIT_ERR_bits); struct DetIdGPU { - uint32_t RawId; + uint32_t rawId; uint32_t rocInDet; uint32_t moduleId; }; @@ -207,19 +207,6 @@ namespace pixelgpudetails { SiPixelDigiErrorsCUDA digiErrors_d; }; - // see RecoLocalTracker/SiPixelClusterizer - // all are runtime const, should be specified in python _cfg.py - struct ADCThreshold { - const int thePixelThreshold = 1000; // default Pixel threshold in electrons - const int theSeedThreshold = 1000; // seed thershold in electrons not used in our algo - const float theClusterThreshold = 4000; // cluster threshold in electron - const int ConversionFactor = 65; // adc to electron conversion factor - - const int theStackADC_ = 255; // the maximum adc count for stack layer - const int theFirstStack_ = 5; // the index of the fits stack layer - const double theElectronPerADCGain_ = 600; // ADC to electron conversion - }; - } // namespace pixelgpudetails #endif // RecoLocalTracker_SiPixelClusterizer_plugins_SiPixelRawToClusterGPUKernel_h diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h index 56bbab083483d..9f295981ca732 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h @@ -36,7 +36,6 @@ namespace gpuClustering { } __global__ - // __launch_bounds__(256,4) void findClus(uint16_t const* __restrict__ id, // module id of each pixel uint16_t const* __restrict__ x, // local coordinates of each pixel diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc index 774727cf91d9c..3dfbc5949422a 100644 --- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc +++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc @@ -309,9 +309,9 @@ void PixelCPEFast::errorFromTemplates(DetParam const& theDetParam, float dummy; SiPixelGenError gtempl(thePixelGenError_); - int gtemplID_ = theDetParam.detTemplateId; + int gtemplID = theDetParam.detTemplateId; - theClusterParam.qBin_ = gtempl.qbin(gtemplID_, + theClusterParam.qBin_ = gtempl.qbin(gtemplID, theClusterParam.cotalpha, theClusterParam.cotbeta, locBz, From 85d69f9d1f247ae554556dfd2ed22f0e96761393 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 08:08:49 +0100 Subject: [PATCH 22/40] General code cleanup General clean up of the pixel local reconstructon code: - improve comments, remove commented out code and obsolete comments - add notes about the original cpu code - update variable names to follow the coding rules --- .../plugins/SiPixelRawToClusterGPUKernel.cu | 42 ++++++++----------- 1 file changed, 18 insertions(+), 24 deletions(-) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 32143e7e5c0a2..25e5c925990f8 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -3,7 +3,6 @@ * File Name: RawToClusterGPU.cu * Description: It converts Raw data into Digi Format on GPU * Finaly the Output of RawToDigi data is given to pixelClusterizer - * **/ // C++ includes @@ -61,7 +60,9 @@ namespace pixelgpudetails { __device__ uint32_t getADC(uint32_t ww) { return ((ww >> pixelgpudetails::ADC_shift) & pixelgpudetails::ADC_mask); } - __device__ bool isBarrel(uint32_t rawId) { return (PixelSubdetector::PixelBarrel == ((rawId >> DetId::kSubdetOffset) & DetId::kSubdetMask)); } + __device__ bool isBarrel(uint32_t rawId) { + return (PixelSubdetector::PixelBarrel == ((rawId >> DetId::kSubdetOffset) & DetId::kSubdetMask)); + } __device__ pixelgpudetails::DetIdGPU getRawId(const SiPixelROCsStatusAndMapping *cablingMap, uint8_t fed, @@ -144,6 +145,7 @@ namespace pixelgpudetails { return global; } + // error decoding and handling copied from EventFilter/SiPixelRawToDigi/src/ErrorChecker.cc __device__ uint8_t conversionError(uint8_t fedId, uint8_t status, bool debug = false) { uint8_t errorType = 0; @@ -187,6 +189,7 @@ namespace pixelgpudetails { __device__ bool dcolIsValid(uint32_t dcol, uint32_t pxid) { return ((dcol < 26) & (2 <= pxid) & (pxid < 162)); } + // error decoding and handling copied from EventFilter/SiPixelRawToDigi/src/ErrorChecker.cc __device__ uint8_t checkROC(uint32_t errorWord, uint8_t fedId, uint32_t link, @@ -240,15 +243,15 @@ namespace pixelgpudetails { case (30): { if (debug) printf("TBM error trailer (errorType = 30)\n"); - int StateMatch_bits = 4; - int StateMatch_shift = 8; - uint32_t StateMatch_mask = ~(~uint32_t(0) << StateMatch_bits); - int StateMatch = (errorWord >> StateMatch_shift) & StateMatch_mask; - if (StateMatch != 1 && StateMatch != 8) { + int stateMatch_bits = 4; + int stateMatch_shift = 8; + uint32_t stateMatch_mask = ~(~uint32_t(0) << stateMatch_bits); + int stateMatch = (errorWord >> stateMatch_shift) & stateMatch_mask; + if (stateMatch != 1 && stateMatch != 8) { if (debug) printf("FED error 30 with unexpected State Bits (errorType = 30)\n"); } - if (StateMatch == 1) + if (stateMatch == 1) errorType = 40; // 1=Overflow -> 40, 8=number of ROCs -> 30 errorFound = true; break; @@ -266,6 +269,7 @@ namespace pixelgpudetails { return errorFound ? errorType : 0; } + // error decoding and handling copied from EventFilter/SiPixelRawToDigi/src/ErrorChecker.cc __device__ uint32_t getErrRawID(uint8_t fedId, uint32_t errWord, uint32_t errorType, @@ -279,13 +283,10 @@ namespace pixelgpudetails { case 31: case 36: case 40: { - //set dummy values for cabling just to get detId from link - //cabling.dcol = 0; - //cabling.pxid = 2; uint32_t roc = 1; uint32_t link = (errWord >> pixelgpudetails::LINK_shift) & pixelgpudetails::LINK_mask; uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; - if (rID_temp != 9999) + if (rID_temp != gpuClustering::invalidModuleId) rID = rID_temp; break; } @@ -315,24 +316,19 @@ namespace pixelgpudetails { if ((chanNmbr < 1) || (chanNmbr > 36)) break; // signifies unexpected result - // set dummy values for cabling just to get detId from link if in Barrel - //cabling.dcol = 0; - //cabling.pxid = 2; uint32_t roc = 1; uint32_t link = chanNmbr; uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; - if (rID_temp != 9999) + if (rID_temp != gpuClustering::invalidModuleId) rID = rID_temp; break; } case 37: case 38: { - //cabling.dcol = 0; - //cabling.pxid = 2; uint32_t roc = (errWord >> pixelgpudetails::ROC_shift) & pixelgpudetails::ROC_mask; uint32_t link = (errWord >> pixelgpudetails::LINK_shift) & pixelgpudetails::LINK_mask; uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; - if (rID_temp != 9999) + if (rID_temp != gpuClustering::invalidModuleId) rID = rID_temp; break; } @@ -374,7 +370,7 @@ namespace pixelgpudetails { // initialize (too many coninue below) pdigi[gIndex] = 0; rawIdArr[gIndex] = 0; - moduleId[gIndex] = 9999; + moduleId[gIndex] = gpuClustering::invalidModuleId; uint32_t ww = word[gIndex]; // Array containing 32 bit raw data if (ww == 0) { @@ -408,8 +404,8 @@ namespace pixelgpudetails { if (skipROC) continue; - uint32_t layer = 0; //, ladder =0; - int side = 0, panel = 0, module = 0; //disk = 0, blade = 0 + uint32_t layer = 0; + int side = 0, panel = 0, module = 0; if (barrel) { layer = (rawId >> pixelgpudetails::layerStartBit) & pixelgpudetails::layerMask; @@ -419,9 +415,7 @@ namespace pixelgpudetails { // endcap ids layer = 0; panel = (rawId >> pixelgpudetails::panelStartBit) & pixelgpudetails::panelMask; - //disk = (rawId >> diskStartBit_) & diskMask_; side = (panel == 1) ? -1 : 1; - //blade = (rawId >> bladeStartBit_) & bladeMask_; } // ***special case of layer to 1 be handled here From 7ae9dde4393cdff2913d83915fff9aae12b2cf67 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 09:46:44 +0100 Subject: [PATCH 23/40] Use std::size instead of hardcoding the array size --- .../interface/phase1PixelTopology.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h index 409ebec3cb43f..c2b5bc9d95f83 100644 --- a/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h +++ b/Geometry/TrackerGeometryBuilder/interface/phase1PixelTopology.h @@ -65,7 +65,7 @@ namespace phase1PixelTopology { bool go = true; int n = 2; while (go) { - for (uint8_t i = 1; i < 11; ++i) { + for (uint8_t i = 1; i < std::size(layerStart); ++i) { if (layerStart[i] % n != 0) { go = false; break; @@ -81,18 +81,18 @@ namespace phase1PixelTopology { constexpr uint32_t maxModuleStride = findMaxModuleStride(); constexpr uint8_t findLayer(uint32_t detId) { - for (uint8_t i = 0; i < 11; ++i) + for (uint8_t i = 0; i < std::size(layerStart); ++i) if (detId < layerStart[i + 1]) return i; - return 11; + return std::size(layerStart); } constexpr uint8_t findLayerFromCompact(uint32_t detId) { detId *= maxModuleStride; - for (uint8_t i = 0; i < 11; ++i) + for (uint8_t i = 0; i < std::size(layerStart); ++i) if (detId < layerStart[i + 1]) return i; - return 11; + return std::size(layerStart); } constexpr uint32_t layerIndexSize = numberOfModules / maxModuleStride; From 949a08a62765a03e93d07cca8981620550250127 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 10:01:17 +0100 Subject: [PATCH 24/40] Update comments --- .../SiPixelClusterizer/plugins/SiPixelClusterProducer.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc index 5a8ea9e772163..2bd902af01b1e 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelClusterProducer.cc @@ -123,7 +123,7 @@ void SiPixelClusterProducer::produce(edm::Event& e, const edm::EventSetup& es) { // Step D: write output to file output->shrink_to_fit(); - // set sequential identifier (this is a const interface, but we need to set it after the sorting) + // set sequential identifier for (auto& clusters : *output) { uint16_t id = 0; for (auto& cluster : clusters) { From 10373579dfa38ad5c9f0524cad4a5f572823e116 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 10:03:02 +0100 Subject: [PATCH 25/40] Move common code to PixelClusterizerBase --- .../plugins/PixelClusterizerBase.h | 38 ++++++++++++------- .../plugins/SiPixelDigisClustersFromSoA.cc | 36 +----------------- 2 files changed, 26 insertions(+), 48 deletions(-) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/PixelClusterizerBase.h b/RecoLocalTracker/SiPixelClusterizer/plugins/PixelClusterizerBase.h index 9e3aad606851c..eb622cccb051e 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/PixelClusterizerBase.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/PixelClusterizerBase.h @@ -1,13 +1,14 @@ #ifndef RecoLocalTracker_SiPixelClusterizer_PixelClusterizerBase_H #define RecoLocalTracker_SiPixelClusterizer_PixelClusterizerBase_H +#include + +#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationServiceBase.h" #include "DataFormats/Common/interface/DetSetVector.h" #include "DataFormats/Common/interface/DetSetVectorNew.h" #include "DataFormats/SiPixelCluster/interface/SiPixelCluster.h" #include "DataFormats/SiPixelDigi/interface/PixelDigi.h" #include "DataFormats/TrackerCommon/interface/TrackerTopology.h" -#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationServiceBase.h" -#include class PixelGeomDetUnit; @@ -20,29 +21,38 @@ class PixelClusterizerBase { typedef edmNew::DetSet::const_iterator ClusterIterator; struct AccretionCluster { - typedef unsigned short UShort; - static constexpr UShort MAXSIZE = 256; - UShort adc[MAXSIZE]; - UShort x[MAXSIZE]; - UShort y[MAXSIZE]; - UShort xmin = 16000; - UShort ymin = 16000; + static constexpr uint16_t MAXSIZE = 256; + uint16_t adc[MAXSIZE]; + uint16_t x[MAXSIZE]; + uint16_t y[MAXSIZE]; + uint16_t xmin = 16000; + uint16_t ymin = 16000; unsigned int isize = 0; - unsigned int curr = 0; + int charge = 0; // stack interface (unsafe ok for use below) - UShort top() const { return curr; } + unsigned int curr = 0; + uint16_t top() const { return curr; } void pop() { ++curr; } bool empty() { return curr == isize; } - bool add(SiPixelCluster::PixelPos const& p, UShort const iadc) { + void clear() { + xmin = 16000; + ymin = 16000; + isize = 0; + charge = 0; + curr = 0; + } + + bool add(SiPixelCluster::PixelPos const& p, uint16_t const iadc) { if (isize == MAXSIZE) return false; - xmin = std::min(xmin, (unsigned short)(p.row())); - ymin = std::min(ymin, (unsigned short)(p.col())); + xmin = std::min(xmin, p.row()); + ymin = std::min(ymin, p.col()); adc[isize] = iadc; x[isize] = p.row(); y[isize++] = p.col(); + charge += iadc; return true; } }; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index af260a46cb8bc..6cd35e41134f5 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -14,39 +14,7 @@ #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/ParameterSet/interface/ParameterSetDescription.h" #include "Geometry/Records/interface/TrackerTopologyRcd.h" - -namespace { - struct AccretionCluster { - typedef unsigned short UShort; - static constexpr UShort MAXSIZE = 256; - UShort adc[MAXSIZE]; - UShort x[MAXSIZE]; - UShort y[MAXSIZE]; - UShort xmin = 16000; - UShort ymin = 16000; - unsigned int isize = 0; - int charge = 0; - - void clear() { - isize = 0; - charge = 0; - xmin = 16000; - ymin = 16000; - } - - bool add(SiPixelCluster::PixelPos const& p, UShort const iadc) { - if (isize == MAXSIZE) - return false; - xmin = std::min(xmin, (unsigned short)(p.row())); - ymin = std::min(ymin, (unsigned short)(p.col())); - adc[isize] = iadc; - x[isize] = p.row(); - y[isize++] = p.col(); - charge += iadc; - return true; - } - }; -} // namespace +#include "RecoLocalTracker/SiPixelClusterizer/plugins/PixelClusterizerBase.h" class SiPixelDigisClustersFromSoA : public edm::global::EDProducer<> { public: @@ -98,7 +66,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con } int32_t nclus = -1; - std::vector aclusters(1024); + std::vector aclusters(1024); auto totCluseFilled = 0; auto fillClusters = [&](uint32_t detId) { From 5ca25bb3836201733749491ad7dc6066c53ce895 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 08:08:49 +0100 Subject: [PATCH 26/40] General code cleanup General clean up of the pixel local reconstructon code: - replace std::cout with LogDebug or LogWarning - reuse existing constants --- .../plugins/SiPixelDigiErrorsFromSoA.cc | 11 +++++----- .../plugins/SiPixelDigisClustersFromSoA.cc | 21 ++++++++++++------- .../plugins/SiPixelRawToClusterCUDA.cc | 2 +- 3 files changed, 21 insertions(+), 13 deletions(-) diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc index f8cb585f604f3..625c4d2c941c0 100644 --- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc +++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc @@ -1,24 +1,25 @@ +#include + #include "CondFormats/DataRecord/interface/SiPixelFedCablingMapRcd.h" #include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingMap.h" #include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingTree.h" #include "DataFormats/Common/interface/DetSetVector.h" #include "DataFormats/Common/interface/Handle.h" #include "DataFormats/DetId/interface/DetIdCollection.h" +#include "DataFormats/FEDRawData/interface/FEDNumbering.h" #include "DataFormats/SiPixelDetId/interface/PixelFEDChannel.h" #include "DataFormats/SiPixelDigi/interface/PixelDigi.h" #include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h" #include "EventFilter/SiPixelRawToDigi/interface/PixelDataFormatter.h" #include "FWCore/Framework/interface/ESTransientHandle.h" #include "FWCore/Framework/interface/ESWatcher.h" -#include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/Event.h" +#include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "FWCore/Framework/interface/stream/EDProducer.h" #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" -#include "FWCore/ParameterSet/interface/ParameterSetDescription.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" - -#include +#include "FWCore/ParameterSet/interface/ParameterSetDescription.h" class SiPixelDigiErrorsFromSoA : public edm::stream::EDProducer<> { public: @@ -98,7 +99,7 @@ void SiPixelDigiErrorsFromSoA::produce(edm::Event& iEvent, const edm::EventSetup for (auto i = 0U; i < size; i++) { SiPixelErrorCompact err = digiErrors.error(i); if (err.errorType != 0) { - SiPixelRawDataError error(err.word, err.errorType, err.fedId + 1200); + SiPixelRawDataError error(err.word, err.errorType, err.fedId + FEDNumbering::MINSiPixeluTCAFEDID); errors[err.rawId].push_back(error); } } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index 6cd35e41134f5..7be36f0a5f963 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -10,6 +10,7 @@ #include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "FWCore/Framework/interface/global/EDProducer.h" +#include "FWCore/MessageLogger/interface/MessageLogger.h" #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/ParameterSet/interface/ParameterSetDescription.h" @@ -66,8 +67,10 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con } int32_t nclus = -1; - std::vector aclusters(1024); - auto totCluseFilled = 0; + std::vector aclusters(gpuClustering::maxNumClustersPerModules); +#ifdef EDM_ML_DEBUG + auto totClustersFilled = 0; +#endif auto fillClusters = [&](uint32_t detId) { if (nclus < 0) @@ -83,8 +86,10 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con << "Layer/DetId/clusId " << layer << '/' << detId << '/' << ic << " size/charge " << acluster.isize << '/' << acluster.charge; SiPixelCluster cluster(acluster.isize, acluster.adc, acluster.x, acluster.y, acluster.xmin, acluster.ymin, ic); - ++totCluseFilled; - // std::cout << "putting in this cluster " << ic << " " << cluster.charge() << " " << cluster.pixelADC().size() << endl; +#ifdef EDM_ML_DEBUG + ++totClustersFilled; +#endif + LogDebug("SiPixelDigisClustersFromSoA") << "putting in this cluster " << ic << " " << cluster.charge() << " " << cluster.pixelADC().size(); // sort by row (x) spc.push_back(std::move(cluster)); std::push_heap(spc.begin(), spc.end(), [](SiPixelCluster const& cl1, SiPixelCluster const& cl2) { @@ -115,14 +120,14 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con if ((*detDigis).empty()) (*detDigis).data.reserve(64); // avoid the first relocations else { - std::cout << "Problem det present twice in input! " << (*detDigis).detId() << std::endl; + edm::LogWarning("SiPixelDigisClustersFromSoA") << "Problem det present twice in input! " << (*detDigis).detId(); } } (*detDigis).data.emplace_back(digis.pdigi(i)); auto const& dig = (*detDigis).data.back(); // fill clusters assert(digis.clus(i) >= 0); - assert(digis.clus(i) < 1024); + assert(digis.clus(i) < gpuClustering::maxNumClustersPerModules); nclus = std::max(digis.clus(i), nclus); auto row = dig.row(); auto col = dig.column(); @@ -133,7 +138,9 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con // fill final clusters if (detDigis) fillClusters((*detDigis).detId()); - //std::cout << "filled " << totCluseFilled << " clusters" << std::endl; +#ifdef EDM_ML_DEBUG + LogDebug("SiPixelDigisClustersFromSoA") << "filled " << totClustersFilled << " clusters"; +#endif iEvent.put(digiPutToken_, std::move(collection)); iEvent.put(clusterPutToken_, std::move(outputClusters)); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc index 5e97610d92286..93b92e145ec5c 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc @@ -184,7 +184,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent, // for GPU // first 150 index stores the fedId and next 150 will store the // start index of word in that fed - assert(fedId >= 1200); + assert(fedId >= FEDNumbering::MINSiPixeluTCAFEDID); fedCounter++; // get event data for this fed From e78b05cd463431617b9266d92cdc646d4a9a4312 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 11:07:10 +0100 Subject: [PATCH 27/40] Apply code formatting --- .../TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc | 3 ++- .../SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc index 7df49b1c9f780..dd3cf548e11dd 100644 --- a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc +++ b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc @@ -14,6 +14,7 @@ cms::cuda::host::unique_ptr TrackingRecHit2DCUDA::localCoordToHostAsync template <> cms::cuda::host::unique_ptr TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cudaStream_t stream) const { auto ret = cms::cuda::make_host_unique(gpuClustering::maxNumModules + 1, stream); - cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClustering::maxNumModules + 1), cudaMemcpyDefault, stream)); + cudaCheck(cudaMemcpyAsync( + ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClustering::maxNumModules + 1), cudaMemcpyDefault, stream)); return ret; } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index 7be36f0a5f963..0078bae38306a 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -89,7 +89,8 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con #ifdef EDM_ML_DEBUG ++totClustersFilled; #endif - LogDebug("SiPixelDigisClustersFromSoA") << "putting in this cluster " << ic << " " << cluster.charge() << " " << cluster.pixelADC().size(); + LogDebug("SiPixelDigisClustersFromSoA") + << "putting in this cluster " << ic << " " << cluster.charge() << " " << cluster.pixelADC().size(); // sort by row (x) spc.push_back(std::move(cluster)); std::push_heap(spc.begin(), spc.end(), [](SiPixelCluster const& cl1, SiPixelCluster const& cl2) { From 7870594ed8850f301ac7b3e8d7c9430328d0916f Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 11:17:47 +0100 Subject: [PATCH 28/40] Convert iterator-based loops to range-based loops --- .../plugins/SiPixelDigiErrorsFromSoA.cc | 6 +++--- .../plugins/SiPixelRecHitFromSOA.cc | 18 +++++++++--------- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 16 ++++++++-------- .../PixelTriplets/plugins/GPUCACell.h | 3 +-- 4 files changed, 21 insertions(+), 22 deletions(-) diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc index 625c4d2c941c0..b487942a1419b 100644 --- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc +++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc @@ -106,13 +106,13 @@ void SiPixelDigiErrorsFromSoA::produce(edm::Event& iEvent, const edm::EventSetup constexpr uint32_t dummydetid = 0xffffffff; typedef PixelDataFormatter::Errors::iterator IE; - for (IE is = errors.begin(); is != errors.end(); is++) { - uint32_t errordetid = is->first; + for (auto& error : errors) { + uint32_t errordetid = error.first; if (errordetid == dummydetid) { // errors given dummy detId must be sorted by Fed nodeterrors.insert(nodeterrors.end(), errors[errordetid].begin(), errors[errordetid].end()); } else { edm::DetSet& errorDetSet = errorcollection.find_or_insert(errordetid); - errorDetSet.data.insert(errorDetSet.data.end(), is->second.begin(), is->second.end()); + errorDetSet.data.insert(errorDetSet.data.end(), error.second.begin(), error.second.end()); // Fill detid of the detectors where there is error AND the error number is listed // in the configurable error list in the job option cfi. // Code needs to be here, because there can be a set of errors for each diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc index b2e7c36cffd10..ea5107c5af90f 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc @@ -111,9 +111,9 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es int numberOfDetUnits = 0; int numberOfClusters = 0; - for (auto DSViter = input.begin(); DSViter != input.end(); DSViter++) { + for (auto const& dsv: input) { numberOfDetUnits++; - unsigned int detid = DSViter->detId(); + unsigned int detid = dsv.detId(); DetId detIdObject(detid); const GeomDetUnit* genericDet = geom->idToDetUnit(detIdObject); auto gind = genericDet->index(); @@ -125,26 +125,26 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es auto nhits = lc - fc; assert(lc > fc); - // std::cout << "in det " << gind << ": conv " << nhits << " hits from " << DSViter->size() << " legacy clusters" + // std::cout << "in det " << gind << ": conv " << nhits << " hits from " << dsv.size() << " legacy clusters" // <<' '<< fc <<','< maxHitsInModule) printf( "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nhits, gind, maxHitsInModule); nhits = std::min(nhits, maxHitsInModule); - //std::cout << "in det " << gind << "conv " << nhits << " hits from " << DSViter->size() << " legacy clusters" + //std::cout << "in det " << gind << "conv " << nhits << " hits from " << dsv.size() << " legacy clusters" // <<' '<< lc <<','<size()); - if (nhits != DSViter->size()) { - edm::LogWarning("GPUHits2CPU") << "nhits!= nclus " << nhits << ' ' << DSViter->size() << std::endl; + assert(nhits <= dsv.size()); + if (nhits != dsv.size()) { + edm::LogWarning("GPUHits2CPU") << "nhits!= nclus " << nhits << ' ' << dsv.size() << std::endl; } - for (auto const& clust : *DSViter) { + for (auto const& clust : dsv) { assert(clust.originalId() >= 0); - assert(clust.originalId() < DSViter->size()); + assert(clust.originalId() < dsv.size()); if (clust.originalId() >= nhits) continue; auto ij = jnd(clust.originalId()); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 26688b1c10003..0ebeb226c66bb 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -125,13 +125,13 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv // fill cluster arrays int numberOfClusters = 0; - for (auto DSViter = input.begin(); DSViter != input.end(); DSViter++) { - unsigned int detid = DSViter->detId(); + for (auto const& dsv : input) { + unsigned int detid = dsv.detId(); DetId detIdObject(detid); const GeomDetUnit* genericDet = geom_->idToDetUnit(detIdObject); auto gind = genericDet->index(); assert(gind < gpuClustering::maxNumModules); - auto const nclus = DSViter->size(); + auto const nclus = dsv.size(); clusInModule_[gind] = nclus; numberOfClusters += nclus; } @@ -155,16 +155,16 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv int numberOfDetUnits = 0; int numberOfHits = 0; - for (auto DSViter = input.begin(); DSViter != input.end(); DSViter++) { + for (auto const& dsv : input) { numberOfDetUnits++; - unsigned int detid = DSViter->detId(); + unsigned int detid = dsv.detId(); DetId detIdObject(detid); const GeomDetUnit* genericDet = geom_->idToDetUnit(detIdObject); auto const gind = genericDet->index(); assert(gind < gpuClustering::maxNumModules); const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); - auto const nclus = DSViter->size(); + auto const nclus = dsv.size(); assert(clusInModule_[gind] == nclus); if (0 == nclus) continue; // is this really possible? @@ -173,7 +173,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv auto const lc = hitsModuleStart[gind + 1]; assert(lc > fc); LogDebug("SiPixelRecHitSoAFromLegacy") << "in det " << gind << ": conv " << nclus << " hits from " - << DSViter->size() << " legacy clusters" << ' ' << fc << ',' << lc; + << dsv.size() << " legacy clusters" << ' ' << fc << ',' << lc; assert((lc - fc) == nclus); if (nclus > maxHitsInModule) printf( @@ -189,7 +189,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv moduleId_ = gind; uint32_t ic = 0; uint32_t ndigi = 0; - for (auto const& clust : *DSViter) { + for (auto const& clust : dsv) { assert(clust.size() > 0); for (int i = 0, nd = clust.size(); i < nd; ++i) { auto px = clust.pixel(i); diff --git a/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h b/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h index ef600489f3e0f..2a74d6a064e73 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/GPUCACell.h @@ -293,8 +293,7 @@ class GPUCACell { assert(tmpNtuplet.size() <= 4); bool last = true; - for (int j = 0; j < outerNeighbors().size(); ++j) { - auto otherCell = outerNeighbors()[j]; + for (unsigned int otherCell : outerNeighbors()) { if (cells[otherCell].theDoubletId < 0) continue; // killed by earlyFishbone last = false; From 6352b74c727d893b0b2ad9a678b83e7f28a32288 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 11:32:45 +0100 Subject: [PATCH 29/40] Convert iterator-based loops to range-based loops --- .../SiPixelRecHits/plugins/SiPixelRecHitConverter.cc | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc index 4ebba652335a0..fabe652747588 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc @@ -193,15 +193,15 @@ namespace cms { auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); auto hitsModuleStart = hmsp.get(); std::array clusInModule{}; - for (auto DSViter = input.begin(); DSViter != input.end(); DSViter++) { - unsigned int detid = DSViter->detId(); + for (auto const& dsv : input) { + unsigned int detid = dsv.detId(); DetId detIdObject(detid); const GeomDetUnit* genericDet = geom.idToDetUnit(detIdObject); auto gind = genericDet->index(); // FIXME to be changed to support Phase2 if (gind >= int(gpuClustering::maxNumModules)) continue; - auto const nclus = DSViter->size(); + auto const nclus = dsv.size(); assert(nclus > 0); clusInModule[gind] = nclus; numberOfClusters += nclus; @@ -216,16 +216,16 @@ namespace cms { iEvent.emplace(tHost_, std::move(hmsp)); // hmsp is gone, hitsModuleStart still alive and kicking... numberOfClusters = 0; - for (auto DSViter = input.begin(); DSViter != input.end(); DSViter++) { + for (auto const& dsv : input) { numberOfDetUnits++; - unsigned int detid = DSViter->detId(); + unsigned int detid = dsv.detId(); DetId detIdObject(detid); const GeomDetUnit* genericDet = geom.idToDetUnit(detIdObject); const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); SiPixelRecHitCollectionNew::FastFiller recHitsOnDetUnit(output, detid); - edmNew::DetSet::const_iterator clustIt = DSViter->begin(), clustEnd = DSViter->end(); + edmNew::DetSet::const_iterator clustIt = dsv.begin(), clustEnd = dsv.end(); for (; clustIt != clustEnd; clustIt++) { numberOfClusters++; From e0fa3916c7aa26a26065d32b0912236b392fb613 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 08:08:49 +0100 Subject: [PATCH 30/40] General code cleanup General clean up of the pixel local reconstructon code: - improve comments - update variable names to follow the coding rules --- .../SiPixelRecHits/interface/pixelCPEforGPU.h | 13 +++++++------ RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc | 3 +-- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h index 681211b82e1af..f655329d02013 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h @@ -128,7 +128,7 @@ namespace pixelCPEforGPU { if (0 == sizeM1) // size 1 return 0; - float W_eff = 0; + float w_eff = 0; bool simple = true; if (1 == sizeM1) { // size 2 //--- Width of the clusters minus the edge (first and last) pixels. @@ -140,14 +140,15 @@ namespace pixelCPEforGPU { auto W_pred = theThickness * cot_angle // geometric correction (in cm) - lorentz_shift; // (in cm) &&& check fpix! - W_eff = std::abs(W_pred) - W_inner; + w_eff = std::abs(W_pred) - W_inner; //--- If the observed charge width is inconsistent with the expectations //--- based on the track, do *not* use W_pred-W_inner. Instead, replace //--- it with an *average* effective charge width, which is the average //--- length of the edge pixels. - simple = - (W_eff < 0.0f) | (W_eff > pitch); // this produces "large" regressions for very small numeric differences... + + // this can produce "large" regressions for very small numeric differences + simple = (w_eff < 0.0f) | (w_eff > pitch); } if (simple) { @@ -157,7 +158,7 @@ namespace pixelCPEforGPU { sum_of_edge += 1.0f; if (last_is_big) sum_of_edge += 1.0f; - W_eff = pitch * 0.5f * sum_of_edge; // ave. length of edge pixels (first+last) (cm) + w_eff = pitch * 0.5f * sum_of_edge; // ave. length of edge pixels (first+last) (cm) } //--- Finally, compute the position in this projection @@ -168,7 +169,7 @@ namespace pixelCPEforGPU { if (Qsum == 0) Qsum = 1.0f; - return 0.5f * (Qdiff / Qsum) * W_eff; + return 0.5f * (Qdiff / Qsum) * w_eff; } constexpr inline void position(CommonParams const& __restrict__ comParams, diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc index 3dfbc5949422a..9b20eb5ffbc1f 100644 --- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc +++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc @@ -1,4 +1,3 @@ -#include #include #include "CondFormats/SiPixelTransient/interface/SiPixelTemplate.h" @@ -44,7 +43,7 @@ PixelCPEFast::PixelCPEFast(edm::ParameterSet const& conf, } // Rechit errors in case other, more correct, errors are not vailable - // This are constants. Maybe there is a more efficienct way to store them. + // These are constants. Maybe there is a more efficienct way to store them. xerr_barrel_l1_ = {0.00115, 0.00120, 0.00088}; xerr_barrel_l1_def_ = 0.01030; yerr_barrel_l1_ = {0.00375, 0.00230, 0.00250, 0.00250, 0.00230, 0.00230, 0.00210, 0.00210, 0.00240}; From e5c83180db87ad2fcf6253130620e20161100125 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 12:16:33 +0100 Subject: [PATCH 31/40] Clarify comments and types regarding HostProduct --- .../plugins/SiPixelRecHitConverter.cc | 11 +++++++---- .../plugins/SiPixelRecHitFromSOA.cc | 16 +++++++++------- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 2 +- .../plugins/PixelTrackProducerFromSoA.cc | 2 +- 4 files changed, 18 insertions(+), 13 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc index fabe652747588..8c16be54e5774 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc @@ -85,7 +85,6 @@ // Make heterogeneous framework happy #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "CUDADataFormats/Common/interface/HostProduct.h" -using HMSstorage = HostProduct; using namespace std; @@ -115,6 +114,8 @@ namespace cms { TrackerGeometry const& geom); private: + using HMSstorage = HostProduct; + // TO DO: maybe allow a map of pointers? /// const PixelClusterParameterEstimator * cpe_; // what we got (for now, one ptr to base class) PixelCPEBase const* cpe_ = nullptr; // What we got (for now, one ptr to base class) @@ -189,9 +190,11 @@ namespace cms { const edmNew::DetSetVector& input = *inputhandle; - // fill cluster arrays + // allocate a buffer for the indices of the clusters auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); + // hitsModuleStart is a non-owning pointer to the buffer auto hitsModuleStart = hmsp.get(); + // fill cluster arrays std::array clusInModule{}; for (auto const& dsv : input) { unsigned int detid = dsv.detId(); @@ -212,8 +215,8 @@ namespace cms { hitsModuleStart[i] = hitsModuleStart[i - 1] + clusInModule[i - 1]; assert(numberOfClusters == int(hitsModuleStart[gpuClustering::maxNumModules])); - // yes a unique ptr of a unique ptr so edm is happy and the pointer stay still... - iEvent.emplace(tHost_, std::move(hmsp)); // hmsp is gone, hitsModuleStart still alive and kicking... + // wrap the buffer in a HostProduct, and move it to the Event, without reallocating the buffer or affecting hitsModuleStart + iEvent.emplace(tHost_, std::move(hmsp)); numberOfClusters = 0; for (auto const& dsv : input) { diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc index ea5107c5af90f..bfe0edd487b8e 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc @@ -28,7 +28,7 @@ class SiPixelRecHitFromSOA : public edm::stream::EDProducer { static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); - using HMSstorage = HostProduct; + using HMSstorage = HostProduct; private: void acquire(edm::Event const& iEvent, @@ -82,12 +82,14 @@ void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, } void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es) { - // yes a unique ptr of a unique ptr so edm is happy - auto sizeOfHitModuleStart = gpuClustering::maxNumModules + 1; - auto hmsp = std::make_unique(sizeOfHitModuleStart); - std::copy(m_hitsModuleStart.get(), m_hitsModuleStart.get() + sizeOfHitModuleStart, hmsp.get()); - auto hms = std::make_unique(std::move(hmsp)); // hmsp is gone - iEvent.put(std::move(hms)); // hms is gone! + + // allocate a buffer for the indices of the clusters + auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); + std::copy(m_hitsModuleStart.get(), m_hitsModuleStart.get() + gpuClustering::maxNumModules + 1, hmsp.get()); + // wrap the buffer in a HostProduct + auto hms = std::make_unique(std::move(hmsp)); + // move the HostProduct to the Event, without reallocating the buffer + iEvent.put(std::move(hms)); auto output = std::make_unique(); if (0 == m_nHits) { diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 0ebeb226c66bb..b58cebcfd7303 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -34,7 +34,7 @@ class SiPixelRecHitSoAFromLegacy : public edm::global::EDProducer<> { static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); using HitModuleStart = std::array; - using HMSstorage = HostProduct; + using HMSstorage = HostProduct; private: void produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const override; diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index 005651ab14493..d6c92ea1e6bcc 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc @@ -48,7 +48,7 @@ class PixelTrackProducerFromSoA : public edm::global::EDProducer<> { static void fillDescriptions(edm::ConfigurationDescriptions &descriptions); // using HitModuleStart = std::array; - using HMSstorage = HostProduct; + using HMSstorage = HostProduct; private: void produce(edm::StreamID streamID, edm::Event &iEvent, const edm::EventSetup &iSetup) const override; From deeb33ab725f3afb9f14c00fd76a22bcfa0e6611 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 16 Dec 2020 14:09:41 +0100 Subject: [PATCH 32/40] Rename siPixelRecHitHostSoA to siPixelRecHitSoAFromLegacy --- .../Configuration/python/customizeHLTforPatatrack.py | 4 ++-- .../SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc | 2 +- .../Configuration/python/customizePixelTracksSoAonCPU.py | 8 ++++---- .../python/customizePixelOnlyForProfiling.py | 2 +- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/HLTrigger/Configuration/python/customizeHLTforPatatrack.py b/HLTrigger/Configuration/python/customizeHLTforPatatrack.py index 44b8b30aca527..13d516e21975e 100644 --- a/HLTrigger/Configuration/python/customizeHLTforPatatrack.py +++ b/HLTrigger/Configuration/python/customizeHLTforPatatrack.py @@ -237,8 +237,8 @@ def customisePixelTrackReconstruction(process): # referenced in process.HLTRecoPixelTracksTask # cpu only: convert the pixel rechits from legacy to SoA format - from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHostSoA_cfi import siPixelRecHitHostSoA as _siPixelRecHitHostSoA - process.hltSiPixelRecHitSoA = _siPixelRecHitHostSoA.clone( + from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromLegacy_cfi import siPixelRecHitSoAFromLegacy as _siPixelRecHitSoAFromLegacy + process.hltSiPixelRecHitSoA = _siPixelRecHitSoAFromLegacy.clone( src = "hltSiPixelClusters", beamSpot = "hltOnlineBeamSpot", convertToLegacy = True diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index b58cebcfd7303..2397434027fa1 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -67,7 +67,7 @@ void SiPixelRecHitSoAFromLegacy::fillDescriptions(edm::ConfigurationDescriptions desc.add("src", edm::InputTag("siPixelClustersPreSplitting")); desc.add("CPE", "PixelCPEFast"); desc.add("convertToLegacy", false); - descriptions.add("siPixelRecHitHostSoA", desc); + descriptions.addWithDefaultLabel(desc); } void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& es) const { diff --git a/RecoPixelVertexing/Configuration/python/customizePixelTracksSoAonCPU.py b/RecoPixelVertexing/Configuration/python/customizePixelTracksSoAonCPU.py index bcd9dbdc51ea4..909959f2d81be 100644 --- a/RecoPixelVertexing/Configuration/python/customizePixelTracksSoAonCPU.py +++ b/RecoPixelVertexing/Configuration/python/customizePixelTracksSoAonCPU.py @@ -9,8 +9,8 @@ def customizePixelTracksSoAonCPU(process): # ensure the same results when running on GPU (which supports only the 'HLT' payload) and CPU process.siPixelClustersPreSplitting.cpu.payloadType = cms.string('HLT') - from RecoLocalTracker.SiPixelRecHits.siPixelRecHitHostSoA_cfi import siPixelRecHitHostSoA - process.siPixelRecHitsPreSplitting = siPixelRecHitHostSoA.clone( + from RecoLocalTracker.SiPixelRecHits.siPixelRecHitSoAFromLegacy_cfi import siPixelRecHitSoAFromLegacy + process.siPixelRecHitsPreSplitting = siPixelRecHitSoAFromLegacy.clone( convertToLegacy = True ) @@ -54,8 +54,8 @@ def customizePixelTracksSoAonCPUForProfiling(process): process.MessageLogger.cerr.FwkReport.reportEvery = 100 process = customizePixelTracksSoAonCPU(process) - process.siPixelRecHitHostSoA.convertToLegacy = False + process.siPixelRecHitSoAFromLegacy.convertToLegacy = False - process.TkSoA = cms.Path(process.offlineBeamSpot+process.siPixelDigis+process.siPixelClustersPreSplitting+process.siPixelRecHitHostSoA+process.pixelTrackSoA+process.pixelVertexSoA) + process.TkSoA = cms.Path(process.offlineBeamSpot + process.siPixelDigis + process.siPixelClustersPreSplitting + process.siPixelRecHitSoAFromLegacy + process.pixelTrackSoA + process.pixelVertexSoA) process.schedule = cms.Schedule(process.TkSoA) return process diff --git a/RecoTracker/Configuration/python/customizePixelOnlyForProfiling.py b/RecoTracker/Configuration/python/customizePixelOnlyForProfiling.py index d46764dbd7edd..24774bbda649c 100644 --- a/RecoTracker/Configuration/python/customizePixelOnlyForProfiling.py +++ b/RecoTracker/Configuration/python/customizePixelOnlyForProfiling.py @@ -25,7 +25,7 @@ def customizePixelOnlyForProfilingGPUOnly(process): # tracks and vertices on the CPU in SoA format, without conversion to legacy format. def customizePixelOnlyForProfilingGPUWithHostCopy(process): - #? process.siPixelRecHitHostSoA.convertToLegacy = False + #? process.siPixelRecHitSoAFromLegacy.convertToLegacy = False process.consumer = cms.EDAnalyzer("GenericConsumer", eventProducts = cms.untracked.vstring('pixelTrackSoA', 'pixelVertexSoA') From 52480b6eb1678187613d4cd7c2d254bf858119a3 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 16:18:24 +0100 Subject: [PATCH 33/40] Rename siPixelClustersCUDAPreSplitting to siPixelClustersPreSplittingCUDA --- EventFilter/SiPixelRawToDigi/python/siPixelDigis_cff.py | 4 ++-- .../python/siPixelClustersPreSplitting_cff.py | 6 +++--- .../SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc | 2 +- .../plugins/ClusterTPAssociationProducerCUDA.cc | 2 +- 4 files changed, 7 insertions(+), 7 deletions(-) diff --git a/EventFilter/SiPixelRawToDigi/python/siPixelDigis_cff.py b/EventFilter/SiPixelRawToDigi/python/siPixelDigis_cff.py index 31ba8596bddc6..5c1ff74be9c69 100644 --- a/EventFilter/SiPixelRawToDigi/python/siPixelDigis_cff.py +++ b/EventFilter/SiPixelRawToDigi/python/siPixelDigis_cff.py @@ -8,10 +8,10 @@ siPixelDigisTask = cms.Task(siPixelDigis) siPixelDigisSoA = _siPixelDigisSoAFromCUDA.clone( - src = "siPixelClustersCUDAPreSplitting" + src = "siPixelClustersPreSplittingCUDA" ) siPixelDigiErrorsSoA = _siPixelDigiErrorsSoAFromCUDA.clone( - src = "siPixelClustersCUDAPreSplitting" + src = "siPixelClustersPreSplittingCUDA" ) siPixelDigiErrors = _siPixelDigiErrorsFromSoA.clone() diff --git a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py index 6839e4582bb2b..8bbf47e9ebf90 100644 --- a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py +++ b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py @@ -8,15 +8,15 @@ siPixelClustersPreSplittingTask = cms.Task(siPixelClustersPreSplitting) -siPixelClustersCUDAPreSplitting = _siPixelRawToClusterCUDA.clone() +siPixelClustersPreSplittingCUDA = _siPixelRawToClusterCUDA.clone() from Configuration.Eras.Modifier_run3_common_cff import run3_common -run3_common.toModify(siPixelClustersCUDAPreSplitting, +run3_common.toModify(siPixelClustersPreSplittingCUDA, isRun2=False ) siPixelDigisClustersPreSplitting = _siPixelDigisClustersFromSoA.clone() siPixelClustersPreSplittingTaskCUDA = cms.Task( - siPixelClustersCUDAPreSplitting, + siPixelClustersPreSplittingCUDA, siPixelDigisClustersPreSplitting, ) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc index ed651db3f71bf..09b90526bf7db 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc @@ -53,7 +53,7 @@ void SiPixelRecHitCUDA::fillDescriptions(edm::ConfigurationDescriptions& descrip edm::ParameterSetDescription desc; desc.add("beamSpot", edm::InputTag("offlineBeamSpotCUDA")); - desc.add("src", edm::InputTag("siPixelClustersCUDAPreSplitting")); + desc.add("src", edm::InputTag("siPixelClustersPreSplittingCUDA")); desc.add("CPE", "PixelCPEFast"); descriptions.add("siPixelRecHitCUDA", desc); } diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc b/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc index 51de45237d639..cd63afa6c7440 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc +++ b/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc @@ -112,7 +112,7 @@ 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("heterogeneousPixelDigiClusterSrc", edm::InputTag("siPixelClustersPreSplittingCUDA")); desc.add("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsCUDAPreSplitting")); desc.add("dumpCSV", false); From dc20554b972d0180339acf6f4f6a8536f0c81d1c Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 16:27:43 +0100 Subject: [PATCH 34/40] Rename siPixelRecHitsCUDAPreSplitting and siPixelRecHitsLegacyPreSplitting to siPixelRecHitsPreSplittingCUDA and siPixelRecHitsPreSplittingLegacy --- .../SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc | 2 +- .../SiPixelRecHits/python/SiPixelRecHits_cfi.py | 8 ++++---- .../plugins/PixelTrackProducerFromSoA.cc | 2 +- .../PixelTriplets/plugins/CAHitNtupletCUDA.cc | 2 +- .../plugins/ClusterTPAssociationProducerCUDA.cc | 2 +- 5 files changed, 8 insertions(+), 8 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc index bfe0edd487b8e..f486e3ef3c715 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc @@ -58,7 +58,7 @@ SiPixelRecHitFromSOA::SiPixelRecHitFromSOA(const edm::ParameterSet& iConfig) void SiPixelRecHitFromSOA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; - desc.add("pixelRecHitSrc", edm::InputTag("siPixelRecHitsCUDAPreSplitting")); + desc.add("pixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA")); desc.add("src", edm::InputTag("siPixelClustersPreSplitting")); descriptions.add("siPixelRecHitFromSOA", desc); } diff --git a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py index 2a0c005e51622..cc81d2cbdfe66 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py +++ b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py @@ -28,14 +28,14 @@ siPixelRecHitsPreSplittingTask = cms.Task(siPixelRecHitsPreSplitting) -siPixelRecHitsCUDAPreSplitting = _siPixelRecHitCUDA.clone( +siPixelRecHitsPreSplittingCUDA = _siPixelRecHitCUDA.clone( beamSpot = "offlineBeamSpotToCUDA" ) -siPixelRecHitsLegacyPreSplitting = _siPixelRecHitFromSOA.clone() +siPixelRecHitsPreSplittingLegacy = _siPixelRecHitFromSOA.clone() siPixelRecHitsPreSplittingTaskCUDA = cms.Task( - siPixelRecHitsCUDAPreSplitting, - siPixelRecHitsLegacyPreSplitting, + siPixelRecHitsPreSplittingCUDA, + siPixelRecHitsPreSplittingLegacy, ) from Configuration.ProcessModifiers.gpu_cff import gpu diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index d6c92ea1e6bcc..cdea22c3a8a24 100644 --- a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc +++ b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc @@ -77,7 +77,7 @@ void PixelTrackProducerFromSoA::fillDescriptions(edm::ConfigurationDescriptions edm::ParameterSetDescription desc; desc.add("beamSpot", edm::InputTag("offlineBeamSpot")); desc.add("trackSrc", edm::InputTag("pixelTrackSoA")); - desc.add("pixelRecHitLegacySrc", edm::InputTag("siPixelRecHitsLegacyPreSplitting")); + desc.add("pixelRecHitLegacySrc", edm::InputTag("siPixelRecHitsPreSplittingLegacy")); desc.add("minNumberOfHits", 0); descriptions.addWithDefaultLabel(desc); diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc index 3f6ea5f43c6f9..3b1ea6fe158b2 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletCUDA.cc @@ -58,7 +58,7 @@ void CAHitNtupletCUDA::fillDescriptions(edm::ConfigurationDescriptions& descript edm::ParameterSetDescription desc; desc.add("onGPU", true); - desc.add("pixelRecHitSrc", edm::InputTag("siPixelRecHitsCUDAPreSplitting")); + desc.add("pixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA")); CAHitNtupletGeneratorOnGPU::fillDescriptions(desc); auto label = "caHitNtupletCUDA"; diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc b/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc index cd63afa6c7440..35337151eda91 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc +++ b/SimTracker/TrackerHitAssociation/plugins/ClusterTPAssociationProducerCUDA.cc @@ -113,7 +113,7 @@ void ClusterTPAssociationProducerCUDA::fillDescriptions(edm::ConfigurationDescri desc.add("phase2OTClusterSrc", edm::InputTag("siPhase2Clusters")); desc.add("trackingParticleSrc", edm::InputTag("mix", "MergedTrackTruth")); desc.add("heterogeneousPixelDigiClusterSrc", edm::InputTag("siPixelClustersPreSplittingCUDA")); - desc.add("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsCUDAPreSplitting")); + desc.add("heterogeneousPixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA")); desc.add("dumpCSV", false); From a12321f7ab1da5dd1d3b34043c660db4809aafb1 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 16:51:45 +0100 Subject: [PATCH 35/40] Minor cleanup of gpuPixelRecHits.h Adjust code after rearranging an #ifdef GPU_DEBUG block. Replace an if check with an assert. --- .../SiPixelRecHits/plugins/gpuPixelRecHits.h | 21 ++++++------------- 1 file changed, 6 insertions(+), 15 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index 1db62e6ead85a..89a40c8723ae3 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -70,7 +70,7 @@ namespace gpuPixelRecHits { #ifdef GPU_DEBUG if (threadIdx.x == 0) { - auto k = first; + auto k = clusters.moduleStart(1 + blockIdx.x); while (digis.moduleInd(k) == invalidModuleId) ++k; assert(digis.moduleInd(k) == me); @@ -84,8 +84,6 @@ namespace gpuPixelRecHits { #endif for (int startClus = 0, endClus = nclus; startClus < endClus; startClus += MaxHitsInIter) { - auto first = clusters.moduleStart(1 + blockIdx.x); - int nClusInIter = std::min(MaxHitsInIter, endClus - startClus); int lastClus = startClus + nClusInIter; assert(nClusInIter <= nclus); @@ -107,12 +105,10 @@ namespace gpuPixelRecHits { clusParams.Q_l_Y[ic] = 0; } - first += threadIdx.x; - __syncthreads(); - // one thead per "digi" - + // one thread per "digi" + auto first = clusters.moduleStart(1 + blockIdx.x) + threadIdx.x; for (int i = first; i < numElements; i += blockDim.x) { auto id = digis.moduleInd(i); if (id == invalidModuleId) @@ -122,11 +118,11 @@ namespace gpuPixelRecHits { auto cl = digis.clus(i); if (cl < startClus || cl >= lastClus) continue; - auto x = digis.xx(i); - auto y = digis.yy(i); cl -= startClus; assert(cl >= 0); assert(cl < MaxHitsInIter); + auto x = digis.xx(i); + auto y = digis.yy(i); atomicMin(&clusParams.minRow[cl], x); atomicMax(&clusParams.maxRow[cl], x); atomicMin(&clusParams.minCol[cl], y); @@ -167,13 +163,10 @@ namespace gpuPixelRecHits { // next one cluster per thread... first = clusters.clusModuleStart(me) + startClus; - for (int ic = threadIdx.x; ic < nClusInIter; ic += blockDim.x) { auto h = first + ic; // output index in global memory - // this cannot happen anymore - if (h >= TrackingRecHit2DSOAView::maxHits()) - break; // overflow... + assert(h < TrackingRecHit2DSOAView::maxHits()); assert(h < hits.nHits()); assert(h < clusters.clusModuleStart(me + 1)); @@ -181,9 +174,7 @@ namespace gpuPixelRecHits { pixelCPEforGPU::errorFromDB(cpeParams->commonParams(), cpeParams->detParams(me), clusParams, ic); // store it - hits.charge(h) = clusParams.charge[ic]; - hits.detectorIndex(h) = me; float xl, yl; From 78337cc9a366536b81525795e6d378834cff44a7 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 17:04:25 +0100 Subject: [PATCH 36/40] Use the autogenerated cfi file instead of PixelCPEFast_cfi.py --- .../python/PixelCPEESProducers_cff.py | 2 +- .../SiPixelRecHits/python/PixelCPEFast_cfi.py | 38 ------------------- 2 files changed, 1 insertion(+), 39 deletions(-) delete mode 100644 RecoLocalTracker/SiPixelRecHits/python/PixelCPEFast_cfi.py diff --git a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py index ea9fab563d164..e3879f4d9d34c 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py +++ b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEESProducers_cff.py @@ -10,7 +10,7 @@ # 2. Pixel Generic CPE # from RecoLocalTracker.SiPixelRecHits.PixelCPEGeneric_cfi import * -from RecoLocalTracker.SiPixelRecHits.PixelCPEFast_cfi import * +from RecoLocalTracker.SiPixelRecHits.PixelCPEFastESProducer_cfi import * # # 3. ESProducer for the Magnetic-field dependent template records # diff --git a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEFast_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/PixelCPEFast_cfi.py deleted file mode 100644 index 3e8957530722c..0000000000000 --- a/RecoLocalTracker/SiPixelRecHits/python/PixelCPEFast_cfi.py +++ /dev/null @@ -1,38 +0,0 @@ -import FWCore.ParameterSet.Config as cms - -PixelCPEFastESProducer = cms.ESProducer("PixelCPEFastESProducer", - - ComponentName = cms.string('PixelCPEFast'), - Alpha2Order = cms.bool(True), - - # Edge cluster errors in microns (determined by looking at residual RMS) - EdgeClusterErrorX = cms.double( 50.0 ), - EdgeClusterErrorY = cms.double( 85.0 ), - - # for CPEBase - useLAWidthFromDB = cms.bool(True), - useLAAlignmentOffsets = cms.bool(False), - - # for CPEBase, used only for testing - lAOffset = cms.double(0), - lAWidthBPix = cms.double(0), - lAWidthFPix = cms.double(0), - - # only for Templates, compute the Lorentz shifts - DoLorentz = cms.bool(False), - - # Can use errors predicted by the template code - # If UseErrorsFromTemplates is False, must also set - # TruncatePixelCharge and LoadTemplatesFromDB to be False - UseErrorsFromTemplates = cms.bool(True), - LoadTemplatesFromDB = cms.bool(True), - - # When set True this gives a slight improvement in resolution at no cost - TruncatePixelCharge = cms.bool(True), - - # petar, for clusterProbability() from TTRHs - ClusterProbComputationFlag = cms.int32(0), - - #MagneticFieldRecord: e.g. "" or "ParabolicMF" - MagneticFieldRecord = cms.ESInputTag(""), -) From 28b7aee2062088008ec7cb09934fb7a25d00c6a9 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 17:36:53 +0100 Subject: [PATCH 37/40] Minor cleanup of SiPixelRecHitFromSOA - remove unused data members; - rename data members for consistency; - migrate to LogDebug and LogWarning; - other minor changes; - apply code formatting. --- .../SiPixelRecHits/plugins/BuildFile.xml | 1 + .../plugins/SiPixelRecHitFromSOA.cc | 85 +++++++++---------- 2 files changed, 39 insertions(+), 47 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml b/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml index 40acdaf2385cb..4457b02203e66 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml +++ b/RecoLocalTracker/SiPixelRecHits/plugins/BuildFile.xml @@ -1,4 +1,5 @@ + diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc index f486e3ef3c715..05c09f99dcdb9 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc @@ -1,7 +1,9 @@ #include -#include "CUDADataFormats/Common/interface/Product.h" +#include + #include "CUDADataFormats/Common/interface/HostProduct.h" +#include "CUDADataFormats/Common/interface/Product.h" #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h" #include "DataFormats/Common/interface/DetSetVectorNew.h" #include "DataFormats/Common/interface/Handle.h" @@ -11,6 +13,7 @@ #include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/MakerMacros.h" #include "FWCore/Framework/interface/stream/EDProducer.h" +#include "FWCore/MessageLogger/interface/MessageLogger.h" #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "FWCore/ParameterSet/interface/ParameterSetDescription.h" @@ -41,10 +44,9 @@ class SiPixelRecHitFromSOA : public edm::stream::EDProducer { edm::EDGetTokenT> tokenHit_; // CUDA hits edm::EDGetTokenT clusterToken_; // Legacy Clusters - uint32_t m_nHits; - cms::cuda::host::unique_ptr m_store16; - cms::cuda::host::unique_ptr m_store32; - cms::cuda::host::unique_ptr m_hitsModuleStart; + uint32_t nHits_; + cms::cuda::host::unique_ptr store32_; + cms::cuda::host::unique_ptr hitsModuleStart_; }; SiPixelRecHitFromSOA::SiPixelRecHitFromSOA(const edm::ParameterSet& iConfig) @@ -60,7 +62,7 @@ void SiPixelRecHitFromSOA::fillDescriptions(edm::ConfigurationDescriptions& desc edm::ParameterSetDescription desc; desc.add("pixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA")); desc.add("src", edm::InputTag("siPixelClustersPreSplitting")); - descriptions.add("siPixelRecHitFromSOA", desc); + descriptions.addWithDefaultLabel(desc); } void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, @@ -70,50 +72,46 @@ void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, cms::cuda::ScopedContextAcquire ctx{inputDataWrapped, std::move(waitingTaskHolder)}; auto const& inputData = ctx.get(inputDataWrapped); - m_nHits = inputData.nHits(); + nHits_ = inputData.nHits(); - // std::cout<< "converting " << m_nHits << " Hits"<< std::endl; + LogDebug("SiPixelRecHitFromSOA") << "converting " << nHits_ << " Hits"; - if (0 == m_nHits) + if (0 == nHits_) return; - m_store32 = inputData.localCoordToHostAsync(ctx.stream()); - // m_store16 = inputData.detIndexToHostAsync(ctx.stream(); - m_hitsModuleStart = inputData.hitsModuleStartToHostAsync(ctx.stream()); + store32_ = inputData.localCoordToHostAsync(ctx.stream()); + hitsModuleStart_ = inputData.hitsModuleStartToHostAsync(ctx.stream()); } void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es) { - // allocate a buffer for the indices of the clusters auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); - std::copy(m_hitsModuleStart.get(), m_hitsModuleStart.get() + gpuClustering::maxNumModules + 1, hmsp.get()); + std::copy(hitsModuleStart_.get(), hitsModuleStart_.get() + gpuClustering::maxNumModules + 1, hmsp.get()); // wrap the buffer in a HostProduct auto hms = std::make_unique(std::move(hmsp)); // move the HostProduct to the Event, without reallocating the buffer iEvent.put(std::move(hms)); auto output = std::make_unique(); - if (0 == m_nHits) { + if (0 == nHits_) { iEvent.put(std::move(output)); return; } - auto xl = m_store32.get(); - auto yl = xl + m_nHits; - auto xe = yl + m_nHits; - auto ye = xe + m_nHits; + auto xl = store32_.get(); + auto yl = xl + nHits_; + auto xe = yl + nHits_; + auto ye = xe + nHits_; const TrackerGeometry* geom = &es.getData(geomToken_); - edm::Handle hclusters; - iEvent.getByToken(clusterToken_, hclusters); - + edm::Handle hclusters = iEvent.getHandle(clusterToken_); auto const& input = *hclusters; constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule(); int numberOfDetUnits = 0; int numberOfClusters = 0; - for (auto const& dsv: input) { + for (auto const& dsv : input) { numberOfDetUnits++; unsigned int detid = dsv.detId(); DetId detIdObject(detid); @@ -122,27 +120,27 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); SiPixelRecHitCollectionNew::FastFiller recHitsOnDetUnit(*output, detid); - auto fc = m_hitsModuleStart[gind]; - auto lc = m_hitsModuleStart[gind + 1]; + auto fc = hitsModuleStart_[gind]; + auto lc = hitsModuleStart_[gind + 1]; auto nhits = lc - fc; assert(lc > fc); - // std::cout << "in det " << gind << ": conv " << nhits << " hits from " << dsv.size() << " legacy clusters" - // <<' '<< fc <<','< maxHitsInModule) - printf( - "WARNING: too many clusters %d in Module %d. Only first %d Hits converted\n", nhits, gind, maxHitsInModule); + edm::LogWarning("SiPixelRecHitFromSOA") << fmt::sprintf( + "Too many clusters %d in module %d. Only the first %d hits will be converted", nhits, gind, maxHitsInModule); nhits = std::min(nhits, maxHitsInModule); - //std::cout << "in det " << gind << "conv " << nhits << " hits from " << dsv.size() << " legacy clusters" - // <<' '<< lc <<','<= 0); @@ -158,11 +156,11 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es numberOfClusters++; - /* cpu version.... (for reference) - std::tuple tuple = cpe_->getParameters( clust, *genericDet ); - LocalPoint lp( std::get<0>(tuple) ); - LocalError le( std::get<1>(tuple) ); - SiPixelRecHitQuality::QualWordType rqw( std::get<2>(tuple) ); + /* cpu version.... (for reference) + std::tuple tuple = cpe_->getParameters( clust, *genericDet ); + LocalPoint lp( std::get<0>(tuple) ); + LocalError le( std::get<1>(tuple) ); + SiPixelRecHitQuality::QualWordType rqw( std::get<2>(tuple) ); */ // Create a persistent edm::Ref to the cluster @@ -174,23 +172,16 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es recHitsOnDetUnit.push_back(hit); // ============================= - // std::cout << "SiPixelRecHitGPUVI " << numberOfClusters << ' '<< lp << " " << le << std::endl; + LogDebug("SiPixelRecHitFromSOA") << "cluster " << numberOfClusters << " at " << lp << ' ' << le; } // <-- End loop on Clusters // LogDebug("SiPixelRecHitGPU") - //std::cout << "SiPixelRecHitGPUVI " - // << " Found " << recHitsOnDetUnit.size() << " RecHits on " << detid //; - // << std::endl; + LogDebug("SiPixelRecHitFromSOA") << "found " << recHitsOnDetUnit.size() << " RecHits on " << detid; } // <-- End loop on DetUnits - /* - std::cout << "SiPixelRecHitGPUVI $ det, clus, lost " - << numberOfDetUnits << ' ' - << numberOfClusters << ' ' - << std::endl; - */ + LogDebug("SiPixelRecHitFromSOA") << "found " << numberOfDetUnits << " dets, " << numberOfClusters << " clusters"; iEvent.put(std::move(output)); } From 6a4e3e5985bcb32b27010346995e3631c50cd40e Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 20:51:26 +0100 Subject: [PATCH 38/40] Rename SiPixelRecHitFromSOA to SiPixelRecHitFromCUDA --- .../python/customizeHLTforPatatrack.py | 4 +-- ...HitFromSOA.cc => SiPixelRecHitFromCUDA.cc} | 36 +++++++++---------- .../python/SiPixelRecHits_cfi.py | 6 ++-- 3 files changed, 23 insertions(+), 23 deletions(-) rename RecoLocalTracker/SiPixelRecHits/plugins/{SiPixelRecHitFromSOA.cc => SiPixelRecHitFromCUDA.cc} (82%) diff --git a/HLTrigger/Configuration/python/customizeHLTforPatatrack.py b/HLTrigger/Configuration/python/customizeHLTforPatatrack.py index 13d516e21975e..a8624de594ec8 100644 --- a/HLTrigger/Configuration/python/customizeHLTforPatatrack.py +++ b/HLTrigger/Configuration/python/customizeHLTforPatatrack.py @@ -185,12 +185,12 @@ def customisePixelLocalReconstruction(process): ) # SwitchProducer wrapping the legacy pixel rechit producer or the transfer of the pixel rechits to the host and the conversion from SoA - from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromSOA_cfi import siPixelRecHitFromSOA as _siPixelRecHitFromSOA + from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromCUDA_cfi import siPixelRecHitFromCUDA as _siPixelRecHitFromCUDA process.hltSiPixelRecHits = SwitchProducerCUDA( # legacy producer cpu = process.hltSiPixelRecHits, # converter to legacy format - cuda = _siPixelRecHitFromSOA.clone( + cuda = _siPixelRecHitFromCUDA.clone( pixelRecHitSrc = "hltSiPixelRecHitsCUDA", src = "hltSiPixelClusters" ) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc similarity index 82% rename from RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc rename to RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc index 05c09f99dcdb9..aede0fac5c2bd 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc @@ -24,10 +24,10 @@ #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h" -class SiPixelRecHitFromSOA : public edm::stream::EDProducer { +class SiPixelRecHitFromCUDA : public edm::stream::EDProducer { public: - explicit SiPixelRecHitFromSOA(const edm::ParameterSet& iConfig); - ~SiPixelRecHitFromSOA() override = default; + explicit SiPixelRecHitFromCUDA(const edm::ParameterSet& iConfig); + ~SiPixelRecHitFromCUDA() override = default; static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); @@ -49,23 +49,23 @@ class SiPixelRecHitFromSOA : public edm::stream::EDProducer { cms::cuda::host::unique_ptr hitsModuleStart_; }; -SiPixelRecHitFromSOA::SiPixelRecHitFromSOA(const edm::ParameterSet& iConfig) +SiPixelRecHitFromCUDA::SiPixelRecHitFromCUDA(const edm::ParameterSet& iConfig) : geomToken_(esConsumes()), tokenHit_( consumes>(iConfig.getParameter("pixelRecHitSrc"))), clusterToken_(consumes(iConfig.getParameter("src"))) { - produces(); + produces(); produces(); } -void SiPixelRecHitFromSOA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { +void SiPixelRecHitFromCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; desc.add("pixelRecHitSrc", edm::InputTag("siPixelRecHitsPreSplittingCUDA")); desc.add("src", edm::InputTag("siPixelClustersPreSplitting")); descriptions.addWithDefaultLabel(desc); } -void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, +void SiPixelRecHitFromCUDA::acquire(edm::Event const& iEvent, edm::EventSetup const& iSetup, edm::WaitingTaskWithArenaHolder waitingTaskHolder) { cms::cuda::Product const& inputDataWrapped = iEvent.get(tokenHit_); @@ -74,7 +74,7 @@ void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, nHits_ = inputData.nHits(); - LogDebug("SiPixelRecHitFromSOA") << "converting " << nHits_ << " Hits"; + LogDebug("SiPixelRecHitFromCUDA") << "converting " << nHits_ << " Hits"; if (0 == nHits_) return; @@ -82,7 +82,7 @@ void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent, hitsModuleStart_ = inputData.hitsModuleStartToHostAsync(ctx.stream()); } -void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es) { +void SiPixelRecHitFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& es) { // allocate a buffer for the indices of the clusters auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); std::copy(hitsModuleStart_.get(), hitsModuleStart_.get() + gpuClustering::maxNumModules + 1, hmsp.get()); @@ -91,7 +91,7 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es // move the HostProduct to the Event, without reallocating the buffer iEvent.put(std::move(hms)); - auto output = std::make_unique(); + auto output = std::make_unique(); if (0 == nHits_) { iEvent.put(std::move(output)); return; @@ -119,20 +119,20 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es auto gind = genericDet->index(); const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); - SiPixelRecHitCollectionNew::FastFiller recHitsOnDetUnit(*output, detid); + SiPixelRecHitCollection::FastFiller recHitsOnDetUnit(*output, detid); auto fc = hitsModuleStart_[gind]; auto lc = hitsModuleStart_[gind + 1]; auto nhits = lc - fc; assert(lc > fc); - LogDebug("SiPixelRecHitFromSOA") << "in det " << gind << ": conv " << nhits << " hits from " << dsv.size() + LogDebug("SiPixelRecHitFromCUDA") << "in det " << gind << ": conv " << nhits << " hits from " << dsv.size() << " legacy clusters" << ' ' << fc << ',' << lc; if (nhits > maxHitsInModule) - edm::LogWarning("SiPixelRecHitFromSOA") << fmt::sprintf( + edm::LogWarning("SiPixelRecHitFromCUDA") << fmt::sprintf( "Too many clusters %d in module %d. Only the first %d hits will be converted", nhits, gind, maxHitsInModule); nhits = std::min(nhits, maxHitsInModule); - LogDebug("SiPixelRecHitFromSOA") << "in det " << gind << "conv " << nhits << " hits from " << dsv.size() + LogDebug("SiPixelRecHitFromCUDA") << "in det " << gind << "conv " << nhits << " hits from " << dsv.size() << " legacy clusters" << ' ' << lc << ',' << fc; if (0 == nhits) @@ -172,18 +172,18 @@ void SiPixelRecHitFromSOA::produce(edm::Event& iEvent, edm::EventSetup const& es recHitsOnDetUnit.push_back(hit); // ============================= - LogDebug("SiPixelRecHitFromSOA") << "cluster " << numberOfClusters << " at " << lp << ' ' << le; + LogDebug("SiPixelRecHitFromCUDA") << "cluster " << numberOfClusters << " at " << lp << ' ' << le; } // <-- End loop on Clusters // LogDebug("SiPixelRecHitGPU") - LogDebug("SiPixelRecHitFromSOA") << "found " << recHitsOnDetUnit.size() << " RecHits on " << detid; + LogDebug("SiPixelRecHitFromCUDA") << "found " << recHitsOnDetUnit.size() << " RecHits on " << detid; } // <-- End loop on DetUnits - LogDebug("SiPixelRecHitFromSOA") << "found " << numberOfDetUnits << " dets, " << numberOfClusters << " clusters"; + LogDebug("SiPixelRecHitFromCUDA") << "found " << numberOfDetUnits << " dets, " << numberOfClusters << " clusters"; iEvent.put(std::move(output)); } -DEFINE_FWK_MODULE(SiPixelRecHitFromSOA); +DEFINE_FWK_MODULE(SiPixelRecHitFromCUDA); diff --git a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py index cc81d2cbdfe66..eb9dbad4934cd 100644 --- a/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py +++ b/RecoLocalTracker/SiPixelRecHits/python/SiPixelRecHits_cfi.py @@ -19,10 +19,10 @@ from Configuration.ProcessModifiers.gpu_cff import gpu from RecoLocalTracker.SiPixelRecHits.siPixelRecHitCUDA_cfi import siPixelRecHitCUDA as _siPixelRecHitCUDA -from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromSOA_cfi import siPixelRecHitFromSOA as _siPixelRecHitFromSOA +from RecoLocalTracker.SiPixelRecHits.siPixelRecHitFromCUDA_cfi import siPixelRecHitFromCUDA as _siPixelRecHitFromCUDA gpu.toModify(siPixelRecHitsPreSplitting, - cuda = _siPixelRecHitFromSOA.clone() + cuda = _siPixelRecHitFromCUDA.clone() ) @@ -32,7 +32,7 @@ beamSpot = "offlineBeamSpotToCUDA" ) -siPixelRecHitsPreSplittingLegacy = _siPixelRecHitFromSOA.clone() +siPixelRecHitsPreSplittingLegacy = _siPixelRecHitFromCUDA.clone() siPixelRecHitsPreSplittingTaskCUDA = cms.Task( siPixelRecHitsPreSplittingCUDA, siPixelRecHitsPreSplittingLegacy, From 90d989d3c02bdf6cf143fafe9719adada6b75230 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 21:08:37 +0100 Subject: [PATCH 39/40] Add put tokens and apply code formatting --- .../plugins/SiPixelRecHitFromCUDA.cc | 40 +++++++++---------- 1 file changed, 19 insertions(+), 21 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc index aede0fac5c2bd..790b0da51ecfb 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc @@ -40,9 +40,10 @@ class SiPixelRecHitFromCUDA : public edm::stream::EDProducer void produce(edm::Event& iEvent, edm::EventSetup const& iSetup) override; const edm::ESGetToken geomToken_; - - edm::EDGetTokenT> tokenHit_; // CUDA hits - edm::EDGetTokenT clusterToken_; // Legacy Clusters + const edm::EDGetTokenT> hitsToken_; // CUDA hits + const edm::EDGetTokenT clusterToken_; // legacy clusters + const edm::EDPutTokenT rechitsPutToken_; // legacy rechits + const edm::EDPutTokenT hostPutToken_; uint32_t nHits_; cms::cuda::host::unique_ptr store32_; @@ -51,12 +52,11 @@ class SiPixelRecHitFromCUDA : public edm::stream::EDProducer SiPixelRecHitFromCUDA::SiPixelRecHitFromCUDA(const edm::ParameterSet& iConfig) : geomToken_(esConsumes()), - tokenHit_( + hitsToken_( consumes>(iConfig.getParameter("pixelRecHitSrc"))), - clusterToken_(consumes(iConfig.getParameter("src"))) { - produces(); - produces(); -} + clusterToken_(consumes(iConfig.getParameter("src"))), + rechitsPutToken_(produces()), + hostPutToken_(produces()) {} void SiPixelRecHitFromCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; @@ -66,9 +66,9 @@ void SiPixelRecHitFromCUDA::fillDescriptions(edm::ConfigurationDescriptions& des } void SiPixelRecHitFromCUDA::acquire(edm::Event const& iEvent, - edm::EventSetup const& iSetup, - edm::WaitingTaskWithArenaHolder waitingTaskHolder) { - cms::cuda::Product const& inputDataWrapped = iEvent.get(tokenHit_); + edm::EventSetup const& iSetup, + edm::WaitingTaskWithArenaHolder waitingTaskHolder) { + cms::cuda::Product const& inputDataWrapped = iEvent.get(hitsToken_); cms::cuda::ScopedContextAcquire ctx{inputDataWrapped, std::move(waitingTaskHolder)}; auto const& inputData = ctx.get(inputDataWrapped); @@ -86,14 +86,12 @@ void SiPixelRecHitFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& e // allocate a buffer for the indices of the clusters auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); std::copy(hitsModuleStart_.get(), hitsModuleStart_.get() + gpuClustering::maxNumModules + 1, hmsp.get()); - // wrap the buffer in a HostProduct - auto hms = std::make_unique(std::move(hmsp)); - // move the HostProduct to the Event, without reallocating the buffer - iEvent.put(std::move(hms)); + // wrap the buffer in a HostProduct, and move it to the Event, without reallocating the buffer or affecting hitsModuleStart + iEvent.emplace(hostPutToken_, std::move(hmsp)); - auto output = std::make_unique(); + SiPixelRecHitCollection output; if (0 == nHits_) { - iEvent.put(std::move(output)); + iEvent.emplace(rechitsPutToken_, std::move(output)); return; } @@ -119,21 +117,21 @@ void SiPixelRecHitFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& e auto gind = genericDet->index(); const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); - SiPixelRecHitCollection::FastFiller recHitsOnDetUnit(*output, detid); + SiPixelRecHitCollection::FastFiller recHitsOnDetUnit(output, detid); auto fc = hitsModuleStart_[gind]; auto lc = hitsModuleStart_[gind + 1]; auto nhits = lc - fc; assert(lc > fc); LogDebug("SiPixelRecHitFromCUDA") << "in det " << gind << ": conv " << nhits << " hits from " << dsv.size() - << " legacy clusters" << ' ' << fc << ',' << lc; + << " legacy clusters" << ' ' << fc << ',' << lc; if (nhits > maxHitsInModule) edm::LogWarning("SiPixelRecHitFromCUDA") << fmt::sprintf( "Too many clusters %d in module %d. Only the first %d hits will be converted", nhits, gind, maxHitsInModule); nhits = std::min(nhits, maxHitsInModule); LogDebug("SiPixelRecHitFromCUDA") << "in det " << gind << "conv " << nhits << " hits from " << dsv.size() - << " legacy clusters" << ' ' << lc << ',' << fc; + << " legacy clusters" << ' ' << lc << ',' << fc; if (0 == nhits) continue; @@ -183,7 +181,7 @@ void SiPixelRecHitFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& e LogDebug("SiPixelRecHitFromCUDA") << "found " << numberOfDetUnits << " dets, " << numberOfClusters << " clusters"; - iEvent.put(std::move(output)); + iEvent.emplace(rechitsPutToken_, std::move(output)); } DEFINE_FWK_MODULE(SiPixelRecHitFromCUDA); From 43edd12b5d8e845f0fa9155076435f51d18dba75 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Thu, 17 Dec 2020 21:21:21 +0100 Subject: [PATCH 40/40] Minor cleanup of PixelCPEFast - update variablenames to follow the coding rules - update comments and remove obselete ones --- .../SiPixelRecHits/interface/PixelCPEFast.h | 17 +++--- .../SiPixelRecHits/src/PixelCPEFast.cc | 58 +++++++++---------- 2 files changed, 35 insertions(+), 40 deletions(-) diff --git a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h index 10e10f7654883..7335aa5e2dfdd 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h @@ -66,11 +66,10 @@ class PixelCPEFast final : public PixelCPEBase { int &Q_l_Y, //!< output, Q last in Y bool truncate); - bool UseErrorsFromTemplates_; - bool TruncatePixelCharge_; - - float EdgeClusterErrorX_; - float EdgeClusterErrorY_; + const float edgeClusterErrorX_; + const float edgeClusterErrorY_; + const bool useErrorsFromTemplates_; + const bool truncatePixelCharge_; std::vector xerr_barrel_l1_, yerr_barrel_l1_, xerr_barrel_ln_; std::vector yerr_barrel_ln_, xerr_endcap_, yerr_endcap_; @@ -80,20 +79,18 @@ class PixelCPEFast final : public PixelCPEBase { //--- DB Error Parametrization object, new light templates std::vector thePixelGenError_; - // allocate it with posix malloc to be ocmpatible with cpu wf + // allocate this with posix malloc to be compatible with the cpu workflow std::vector m_detParamsGPU; - // std::vector> m_detParamsGPU; pixelCPEforGPU::CommonParams m_commonParamsGPU; pixelCPEforGPU::LayerGeometry m_layerGeometry; pixelCPEforGPU::AverageGeometry m_averageGeometry; - pixelCPEforGPU::ParamsOnGPU cpuData_; struct GPUData { ~GPUData(); // not needed if not used on CPU... - pixelCPEforGPU::ParamsOnGPU h_paramsOnGPU; - pixelCPEforGPU::ParamsOnGPU *d_paramsOnGPU = nullptr; // copy of the above on the Device + pixelCPEforGPU::ParamsOnGPU paramsOnGPU_h; + pixelCPEforGPU::ParamsOnGPU *paramsOnGPU_d = nullptr; // copy of the above on the Device }; cms::cuda::ESProduct gpuData_; diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc index 9b20eb5ffbc1f..548119cef501b 100644 --- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc +++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc @@ -27,15 +27,13 @@ PixelCPEFast::PixelCPEFast(edm::ParameterSet const& conf, const SiPixelLorentzAngle* lorentzAngle, const SiPixelGenErrorDBObject* genErrorDBObject, const SiPixelLorentzAngle* lorentzAngleWidth) - : PixelCPEBase(conf, mag, geom, ttopo, lorentzAngle, genErrorDBObject, nullptr, lorentzAngleWidth, 0) { - EdgeClusterErrorX_ = conf.getParameter("EdgeClusterErrorX"); - EdgeClusterErrorY_ = conf.getParameter("EdgeClusterErrorY"); - - UseErrorsFromTemplates_ = conf.getParameter("UseErrorsFromTemplates"); - TruncatePixelCharge_ = conf.getParameter("TruncatePixelCharge"); - + : PixelCPEBase(conf, mag, geom, ttopo, lorentzAngle, genErrorDBObject, nullptr, lorentzAngleWidth, 0), + edgeClusterErrorX_(conf.getParameter("EdgeClusterErrorX")), + edgeClusterErrorY_(conf.getParameter("EdgeClusterErrorY")), + useErrorsFromTemplates_(conf.getParameter("UseErrorsFromTemplates")), + truncatePixelCharge_(conf.getParameter("TruncatePixelCharge")) { // Use errors from templates or from GenError - if (UseErrorsFromTemplates_) { + if (useErrorsFromTemplates_) { if (!SiPixelGenError::pushfile(*genErrorDBObject_, thePixelGenError_)) throw cms::Exception("InvalidCalibrationLoaded") << "ERROR: GenErrors not filled correctly. Check the sqlite file. Using SiPixelTemplateDBObject version " @@ -70,37 +68,37 @@ PixelCPEFast::PixelCPEFast(edm::ParameterSet const& conf, const pixelCPEforGPU::ParamsOnGPU* PixelCPEFast::getGPUProductAsync(cudaStream_t cudaStream) const { const auto& data = gpuData_.dataForCurrentDeviceAsync(cudaStream, [this](GPUData& data, cudaStream_t stream) { // and now copy to device... - cudaCheck(cudaMalloc((void**)&data.h_paramsOnGPU.m_commonParams, sizeof(pixelCPEforGPU::CommonParams))); - cudaCheck(cudaMalloc((void**)&data.h_paramsOnGPU.m_detParams, + cudaCheck(cudaMalloc((void**)&data.paramsOnGPU_h.m_commonParams, sizeof(pixelCPEforGPU::CommonParams))); + cudaCheck(cudaMalloc((void**)&data.paramsOnGPU_h.m_detParams, this->m_detParamsGPU.size() * sizeof(pixelCPEforGPU::DetParams))); - cudaCheck(cudaMalloc((void**)&data.h_paramsOnGPU.m_averageGeometry, sizeof(pixelCPEforGPU::AverageGeometry))); - cudaCheck(cudaMalloc((void**)&data.h_paramsOnGPU.m_layerGeometry, sizeof(pixelCPEforGPU::LayerGeometry))); - cudaCheck(cudaMalloc((void**)&data.d_paramsOnGPU, sizeof(pixelCPEforGPU::ParamsOnGPU))); + cudaCheck(cudaMalloc((void**)&data.paramsOnGPU_h.m_averageGeometry, sizeof(pixelCPEforGPU::AverageGeometry))); + cudaCheck(cudaMalloc((void**)&data.paramsOnGPU_h.m_layerGeometry, sizeof(pixelCPEforGPU::LayerGeometry))); + cudaCheck(cudaMalloc((void**)&data.paramsOnGPU_d, sizeof(pixelCPEforGPU::ParamsOnGPU))); cudaCheck(cudaMemcpyAsync( - data.d_paramsOnGPU, &data.h_paramsOnGPU, sizeof(pixelCPEforGPU::ParamsOnGPU), cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync((void*)data.h_paramsOnGPU.m_commonParams, + data.paramsOnGPU_d, &data.paramsOnGPU_h, sizeof(pixelCPEforGPU::ParamsOnGPU), cudaMemcpyDefault, stream)); + cudaCheck(cudaMemcpyAsync((void*)data.paramsOnGPU_h.m_commonParams, &this->m_commonParamsGPU, sizeof(pixelCPEforGPU::CommonParams), cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync((void*)data.h_paramsOnGPU.m_averageGeometry, + cudaCheck(cudaMemcpyAsync((void*)data.paramsOnGPU_h.m_averageGeometry, &this->m_averageGeometry, sizeof(pixelCPEforGPU::AverageGeometry), cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync((void*)data.h_paramsOnGPU.m_layerGeometry, + cudaCheck(cudaMemcpyAsync((void*)data.paramsOnGPU_h.m_layerGeometry, &this->m_layerGeometry, sizeof(pixelCPEforGPU::LayerGeometry), cudaMemcpyDefault, stream)); - cudaCheck(cudaMemcpyAsync((void*)data.h_paramsOnGPU.m_detParams, + cudaCheck(cudaMemcpyAsync((void*)data.paramsOnGPU_h.m_detParams, this->m_detParamsGPU.data(), this->m_detParamsGPU.size() * sizeof(pixelCPEforGPU::DetParams), cudaMemcpyDefault, stream)); }); - return data.d_paramsOnGPU; + return data.paramsOnGPU_d; } void PixelCPEFast::fillParamsForGpu() { @@ -276,12 +274,12 @@ void PixelCPEFast::fillParamsForGpu() { } PixelCPEFast::GPUData::~GPUData() { - if (d_paramsOnGPU != nullptr) { - cudaFree((void*)h_paramsOnGPU.m_commonParams); - cudaFree((void*)h_paramsOnGPU.m_detParams); - cudaFree((void*)h_paramsOnGPU.m_averageGeometry); - cudaFree((void*)h_paramsOnGPU.m_layerGeometry); - cudaFree(d_paramsOnGPU); + if (paramsOnGPU_d != nullptr) { + cudaFree((void*)paramsOnGPU_h.m_commonParams); + cudaFree((void*)paramsOnGPU_h.m_detParams); + cudaFree((void*)paramsOnGPU_h.m_averageGeometry); + cudaFree((void*)paramsOnGPU_h.m_layerGeometry); + cudaFree(paramsOnGPU_d); } } @@ -350,7 +348,7 @@ LocalPoint PixelCPEFast::localPosition(DetParam const& theDetParam, ClusterParam assert(!theClusterParam.with_track_angle); - if (UseErrorsFromTemplates_) { + if (useErrorsFromTemplates_) { errorFromTemplates(theDetParam, theClusterParam, theClusterParam.theCluster->charge()); } else { theClusterParam.qBin_ = 0; @@ -360,7 +358,7 @@ LocalPoint PixelCPEFast::localPosition(DetParam const& theDetParam, ClusterParam int Q_l_X; //!< Q of the last pixel in X int Q_f_Y; //!< Q of the first pixel in Y int Q_l_Y; //!< Q of the last pixel in Y - collect_edge_charges(theClusterParam, Q_f_X, Q_l_X, Q_f_Y, Q_l_Y, UseErrorsFromTemplates_ && TruncatePixelCharge_); + collect_edge_charges(theClusterParam, Q_f_X, Q_l_X, Q_f_Y, Q_l_Y, useErrorsFromTemplates_ && truncatePixelCharge_); // do GPU like ... pixelCPEforGPU::ClusParams cp; @@ -445,8 +443,8 @@ LocalError PixelCPEFast::localError(DetParam const& theDetParam, ClusterParam& t // Default errors are the maximum error used for edge clusters. // These are determined by looking at residuals for edge clusters - float xerr = EdgeClusterErrorX_ * micronsToCm; - float yerr = EdgeClusterErrorY_ * micronsToCm; + float xerr = edgeClusterErrorX_ * micronsToCm; + float yerr = edgeClusterErrorY_ * micronsToCm; // Find if cluster is at the module edge. int maxPixelCol = theClusterParam.theCluster->maxPixelCol(); @@ -464,7 +462,7 @@ LocalError PixelCPEFast::localError(DetParam const& theDetParam, ClusterParam& t bool bigInX = theDetParam.theRecTopol->containsBigPixelInX(minPixelRow, maxPixelRow); bool bigInY = theDetParam.theRecTopol->containsBigPixelInY(minPixelCol, maxPixelCol); - if (UseErrorsFromTemplates_) { + if (useErrorsFromTemplates_) { // // Use template errors