From 5232bcb127d039d99f63e52b2b57e3df17592253 Mon Sep 17 00:00:00 2001 From: Angela Czirkos Date: Tue, 9 Mar 2021 10:29:36 +0100 Subject: [PATCH 1/5] Merge namespaces in gpuClusteringConstants.h --- .../interface/gpuClusteringConstants.h | 16 ++++------------ .../PixelTriplets/plugins/CAConstants.h | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cc | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cu | 2 +- 4 files changed, 7 insertions(+), 15 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index e9dfed7bca7a6..5bd540e69ab44 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -4,31 +4,23 @@ #include #include -namespace pixelGPUConstants { -#ifdef GPU_SMALL_EVENTS - // kept for testing and debugging - constexpr uint32_t maxNumberOfHits = 24 * 1024; -#else - // 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; } + constexpr uint32_t maxNumberOfHits = 24 * 1024; #else // optimized for real data PU 50 // tested on MC events with 55-75 pileup events constexpr uint32_t maxHitsInIter() { return 160; } + // data at pileup 50 has 18300 +/- 3500 hits; 40000 is around 6 sigma away + constexpr uint32_t maxNumberOfHits = 48 * 1024; #endif constexpr uint32_t maxHitsInModule() { return 1024; } constexpr uint16_t maxNumModules = 2000; constexpr int32_t maxNumClustersPerModules = maxHitsInModule(); - constexpr uint32_t maxNumClusters = pixelGPUConstants::maxNumberOfHits; + constexpr uint32_t maxNumClusters = gpuClustering::maxNumberOfHits; constexpr uint16_t invalidModuleId = std::numeric_limits::max() - 1; static_assert(invalidModuleId > maxNumModules); // invalidModuleId must be > maxNumModules diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h index d9c3ff70e35ed..5c31c31ca62f6 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h @@ -57,7 +57,7 @@ namespace CAConstants { using OuterHitOfCell = cms::cuda::VecArray; using TuplesContainer = cms::cuda::OneToManyAssoc; using HitToTuple = - cms::cuda::OneToManyAssoc; // 3.5 should be enough + cms::cuda::OneToManyAssoc; // 3.5 should be enough using TupleMultiplicity = cms::cuda::OneToManyAssoc; } // namespace CAConstants diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc index 1646cb503ff81..a4f032d21cc41 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc @@ -78,7 +78,7 @@ void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA * cms::cuda::launchZero(tuples_d, cudaStream); auto nhits = hh.nHits(); - assert(nhits <= pixelGPUConstants::maxNumberOfHits); + assert(nhits <= gpuClustering::maxNumberOfHits); // std::cout << "N hits " << nhits << std::endl; // if (nhits<2) std::cout << "too few hits " << nhits << std::endl; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu index a8dac7992f4fa..edba7d9cb746a 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu @@ -24,7 +24,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA * cms::cuda::launchZero(tuples_d, cudaStream); auto nhits = hh.nHits(); - assert(nhits <= pixelGPUConstants::maxNumberOfHits); + assert(nhits <= gpuClustering::maxNumberOfHits); // std::cout << "N hits " << nhits << std::endl; // if (nhits<2) std::cout << "too few hits " << nhits << std::endl; From c1be1254d059802a85a941d81b7a2a9bd92da59c Mon Sep 17 00:00:00 2001 From: Angela Czirkos Date: Tue, 9 Mar 2021 11:18:03 +0100 Subject: [PATCH 2/5] Use CUDA Math API functions instead of standard library ones in kernels --- .../plugins/SiPixelRawToClusterGPUKernel.cu | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 25e5c925990f8..97f5f9022f192 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -471,7 +471,8 @@ namespace pixelgpudetails { // limit to maxHitsInModule() for (int i = first, iend = gpuClustering::maxNumModules; i < iend; i += blockDim.x) { - moduleStart[i + 1] = std::min(gpuClustering::maxHitsInModule(), cluStart[i]); + // use CUDA ::min() instead of std::min to allow using constexpr variables in device code + moduleStart[i + 1] = ::min(gpuClustering::maxHitsInModule(), cluStart[i]); } __shared__ uint32_t ws[32]; @@ -485,7 +486,7 @@ namespace pixelgpudetails { #ifdef GPU_DEBUG assert(0 == moduleStart[0]); - auto c0 = std::min(gpuClustering::maxHitsInModule(), cluStart[0]); + auto c0 = ::min(gpuClustering::maxHitsInModule(), cluStart[0]); assert(c0 == moduleStart[1]); assert(moduleStart[1024] >= moduleStart[1023]); assert(moduleStart[1025] >= moduleStart[1024]); @@ -504,7 +505,10 @@ namespace pixelgpudetails { // avoid overflow 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); + // clamp moduleStart[i] to boundaries 0 and maxNumClusters (uint32_t) + // there is not CUDA ::clamp(), replace with ::min() and ::max() + moduleStart[i] = ::max(0U, moduleStart[i]); + moduleStart[i] = ::min(moduleStart[i], maxNumClusters); } } From 1dbc9c7e9aa5d0d3d8b6146dccd329651ad7a198 Mon Sep 17 00:00:00 2001 From: Angela Czirkos Date: Tue, 9 Mar 2021 11:29:23 +0100 Subject: [PATCH 3/5] Use constexpr variables instead of functions for consistency and aesthetics For more info see G.4.16.5. Constexpr variables in https://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf Previous commit was necessary for these changes, since std::min() takes arguments by reference --- .../SiPixelCluster/interface/gpuClusteringConstants.h | 9 ++++----- .../plugins/SiPixelRawToClusterGPUKernel.cu | 6 +++--- .../SiPixelRecHits/interface/pixelCPEforGPU.h | 2 +- .../SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc | 2 +- .../SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc | 2 +- 5 files changed, 10 insertions(+), 11 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index 5bd540e69ab44..260a0c0a8979f 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -7,19 +7,18 @@ namespace gpuClustering { #ifdef GPU_SMALL_EVENTS // kept for testing and debugging - constexpr uint32_t maxHitsInIter() { return 64; } + constexpr uint32_t maxHitsInIter = 64; constexpr uint32_t maxNumberOfHits = 24 * 1024; #else // optimized for real data PU 50 // tested on MC events with 55-75 pileup events - constexpr uint32_t maxHitsInIter() { return 160; } + constexpr uint32_t maxHitsInIter = 160; // data at pileup 50 has 18300 +/- 3500 hits; 40000 is around 6 sigma away constexpr uint32_t maxNumberOfHits = 48 * 1024; #endif - constexpr uint32_t maxHitsInModule() { return 1024; } - + constexpr uint32_t maxHitsInModule = 1024; constexpr uint16_t maxNumModules = 2000; - constexpr int32_t maxNumClustersPerModules = maxHitsInModule(); + constexpr int32_t maxNumClustersPerModules = maxHitsInModule; constexpr uint32_t maxNumClusters = gpuClustering::maxNumberOfHits; constexpr uint16_t invalidModuleId = std::numeric_limits::max() - 1; static_assert(invalidModuleId > maxNumModules); // invalidModuleId must be > maxNumModules diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 97f5f9022f192..88b39aaa26631 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -469,10 +469,10 @@ namespace pixelgpudetails { int first = threadIdx.x; - // limit to maxHitsInModule() + // limit to maxHitsInModule for (int i = first, iend = gpuClustering::maxNumModules; i < iend; i += blockDim.x) { // use CUDA ::min() instead of std::min to allow using constexpr variables in device code - moduleStart[i + 1] = ::min(gpuClustering::maxHitsInModule(), cluStart[i]); + moduleStart[i + 1] = ::min(gpuClustering::maxHitsInModule, cluStart[i]); } __shared__ uint32_t ws[32]; @@ -486,7 +486,7 @@ namespace pixelgpudetails { #ifdef GPU_DEBUG assert(0 == moduleStart[0]); - auto c0 = ::min(gpuClustering::maxHitsInModule(), cluStart[0]); + auto c0 = ::min(gpuClustering::maxHitsInModule, cluStart[0]); assert(c0 == moduleStart[1]); assert(moduleStart[1024] >= moduleStart[1023]); assert(moduleStart[1025] >= moduleStart[1024]); diff --git a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h index 03e136d8d23ef..23f9056cc1d67 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h @@ -98,7 +98,7 @@ namespace pixelCPEforGPU { int16_t ysize[N]; }; - constexpr int32_t MaxHitsInIter = gpuClustering::maxHitsInIter(); + constexpr int32_t MaxHitsInIter = gpuClustering::maxHitsInIter; using ClusParams = ClusParamsT; constexpr inline void computeAnglesFromDet( diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc index 790b0da51ecfb..6b717175c31bc 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc @@ -105,7 +105,7 @@ void SiPixelRecHitFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& e edm::Handle hclusters = iEvent.getHandle(clusterToken_); auto const& input = *hclusters; - constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule(); + constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule; int numberOfDetUnits = 0; int numberOfClusters = 0; diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index 0d95370f25e47..b9bdb9a5fa7a1 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -110,7 +110,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_; From 7b93d9cc3a4df26c9a342aa4dc2c194caf27e452 Mon Sep 17 00:00:00 2001 From: Angela Czirkos Date: Tue, 9 Mar 2021 11:53:38 +0100 Subject: [PATCH 4/5] Change namespace gpuClustering to gpuClusteringConstants --- .../interface/gpuClusteringConstants.h | 6 +-- .../interface/TrackingRecHit2DSOAView.h | 4 +- .../src/TrackingRecHit2DHeterogeneous.cc | 4 +- .../interface/SiPixelGainForHLTonGPU.h | 2 +- .../plugins/SiPixelDigisClustersFromSoA.cc | 6 +-- .../plugins/SiPixelRawToClusterGPUKernel.cu | 38 +++++++++---------- .../plugins/gpuCalibPixel.h | 4 +- .../plugins/gpuClusterChargeCut.h | 4 +- .../plugins/gpuClustering.h | 4 +- .../SiPixelClusterizer/test/gpuClustering_t.h | 2 +- .../SiPixelRecHits/interface/pixelCPEforGPU.h | 2 +- .../plugins/SiPixelRecHitConverter.cc | 10 ++--- .../plugins/SiPixelRecHitFromCUDA.cc | 6 +-- .../plugins/SiPixelRecHitSoAFromLegacy.cc | 20 +++++----- .../SiPixelRecHits/plugins/gpuPixelRecHits.h | 2 +- .../plugins/PixelTrackProducerFromSoA.cc | 2 +- .../PixelTriplets/plugins/CAConstants.h | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cc | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cu | 2 +- .../plugins/gpuPixelDoubletsAlgos.h | 4 +- .../plugins/ClusterSLOnGPU.cu | 4 +- 21 files changed, 65 insertions(+), 65 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h index 260a0c0a8979f..84737de2dd017 100644 --- a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h +++ b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h @@ -4,7 +4,7 @@ #include #include -namespace gpuClustering { +namespace gpuClusteringConstants { #ifdef GPU_SMALL_EVENTS // kept for testing and debugging constexpr uint32_t maxHitsInIter = 64; @@ -19,10 +19,10 @@ namespace gpuClustering { constexpr uint32_t maxHitsInModule = 1024; constexpr uint16_t maxNumModules = 2000; constexpr int32_t maxNumClustersPerModules = maxHitsInModule; - constexpr uint32_t maxNumClusters = gpuClustering::maxNumberOfHits; + constexpr uint32_t maxNumClusters = gpuClusteringConstants::maxNumberOfHits; constexpr uint16_t invalidModuleId = std::numeric_limits::max() - 1; static_assert(invalidModuleId > maxNumModules); // invalidModuleId must be > maxNumModules -} // namespace gpuClustering +} // namespace gpuClusteringConstants #endif // CUDADataFormats_SiPixelCluster_interface_gpuClusteringConstants_h diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h index 7f3c59cd70faf..789f03a22ef98 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 gpuClusteringConstants::maxNumClusters; } using hindex_type = uint32_t; // if above is <=2^32 using PhiBinner = - cms::cuda::HistoContainer; + cms::cuda::HistoContainer; using AverageGeometry = phase1PixelTopology::AverageGeometry; diff --git a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc index dd3cf548e11dd..5785647874558 100644 --- a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc +++ b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc @@ -13,8 +13,8 @@ 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); + auto ret = cms::cuda::make_host_unique(gpuClusteringConstants::maxNumModules + 1, stream); cudaCheck(cudaMemcpyAsync( - ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClustering::maxNumModules + 1), cudaMemcpyDefault, stream)); + ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClusteringConstants::maxNumModules + 1), cudaMemcpyDefault, stream)); return ret; } diff --git a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h index aa5a127927b90..45171ede432ba 100644 --- a/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h +++ b/CondFormats/SiPixelObjects/interface/SiPixelGainForHLTonGPU.h @@ -60,7 +60,7 @@ class SiPixelGainForHLTonGPU { constexpr float decodePed(unsigned int ped) const { return ped * pedPrecision_ + minPed_; } DecodingStructure* v_pedestals_; - std::pair rangeAndCols_[gpuClustering::maxNumModules]; + std::pair rangeAndCols_[gpuClusteringConstants::maxNumModules]; float minPed_, maxPed_, minGain_, maxGain_; float pedPrecision_, gainPrecision_; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index 0078bae38306a..b98d7a8746981 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -54,7 +54,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con auto collection = std::make_unique>(); auto outputClusters = std::make_unique(); - outputClusters->reserve(gpuClustering::maxNumModules, nDigis / 4); + outputClusters->reserve(gpuClusteringConstants::maxNumModules, nDigis / 4); edm::DetSet* detDigis = nullptr; for (uint32_t i = 0; i < nDigis; i++) { @@ -67,7 +67,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con } int32_t nclus = -1; - std::vector aclusters(gpuClustering::maxNumClustersPerModules); + std::vector aclusters(gpuClusteringConstants::maxNumClustersPerModules); #ifdef EDM_ML_DEBUG auto totClustersFilled = 0; #endif @@ -128,7 +128,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con auto const& dig = (*detDigis).data.back(); // fill clusters assert(digis.clus(i) >= 0); - assert(digis.clus(i) < gpuClustering::maxNumClustersPerModules); + assert(digis.clus(i) < gpuClusteringConstants::maxNumClustersPerModules); nclus = std::max(digis.clus(i), nclus); auto row = dig.row(); auto col = dig.column(); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 88b39aaa26631..dc6d3a09a05ff 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -286,7 +286,7 @@ namespace pixelgpudetails { 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 != gpuClustering::invalidModuleId) + if (rID_temp != gpuClusteringConstants::invalidModuleId) rID = rID_temp; break; } @@ -319,7 +319,7 @@ namespace pixelgpudetails { uint32_t roc = 1; uint32_t link = chanNmbr; uint32_t rID_temp = getRawId(cablingMap, fedId, link, roc).rawId; - if (rID_temp != gpuClustering::invalidModuleId) + if (rID_temp != gpuClusteringConstants::invalidModuleId) rID = rID_temp; break; } @@ -328,7 +328,7 @@ namespace pixelgpudetails { 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 != gpuClustering::invalidModuleId) + if (rID_temp != gpuClusteringConstants::invalidModuleId) rID = rID_temp; break; } @@ -370,7 +370,7 @@ namespace pixelgpudetails { // initialize (too many coninue below) pdigi[gIndex] = 0; rawIdArr[gIndex] = 0; - moduleId[gIndex] = gpuClustering::invalidModuleId; + moduleId[gIndex] = gpuClusteringConstants::invalidModuleId; uint32_t ww = word[gIndex]; // Array containing 32 bit raw data if (ww == 0) { @@ -463,48 +463,48 @@ 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(gpuClusteringConstants::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) { + for (int i = first, iend = gpuClusteringConstants::maxNumModules; i < iend; i += blockDim.x) { // use CUDA ::min() instead of std::min to allow using constexpr variables in device code - moduleStart[i + 1] = ::min(gpuClustering::maxHitsInModule, cluStart[i]); + moduleStart[i + 1] = ::min(gpuClusteringConstants::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, gpuClusteringConstants::maxNumModules - 1024, ws); - for (int i = first + 1025, iend = gpuClustering::maxNumModules + 1; i < iend; i += blockDim.x) { + for (int i = first + 1025, iend = gpuClusteringConstants::maxNumModules + 1; i < iend; i += blockDim.x) { moduleStart[i] += moduleStart[1024]; } __syncthreads(); #ifdef GPU_DEBUG assert(0 == moduleStart[0]); - auto c0 = ::min(gpuClustering::maxHitsInModule, cluStart[0]); + auto c0 = ::min(gpuClusteringConstants::maxHitsInModule, cluStart[0]); assert(c0 == moduleStart[1]); assert(moduleStart[1024] >= moduleStart[1023]); assert(moduleStart[1025] >= moduleStart[1024]); - assert(moduleStart[gpuClustering::maxNumModules] >= moduleStart[1025]); + assert(moduleStart[gpuClusteringConstants::maxNumModules] >= moduleStart[1025]); - for (int i = first, iend = gpuClustering::maxNumModules + 1; i < iend; i += blockDim.x) { + for (int i = first, iend = gpuClusteringConstants::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 == gpuClusteringConstants::maxNumModules) printf("moduleStart %d %d\n", i, moduleStart[i]); } #endif // avoid overflow - auto constexpr maxNumClusters = gpuClustering::maxNumClusters; - for (int i = first, iend = gpuClustering::maxNumModules + 1; i < iend; i += blockDim.x) { + auto constexpr maxNumClusters = gpuClusteringConstants::maxNumClusters; + for (int i = first, iend = gpuClusteringConstants::maxNumModules + 1; i < iend; i += blockDim.x) { // clamp moduleStart[i] to boundaries 0 and maxNumClusters (uint32_t) // there is not CUDA ::clamp(), replace with ::min() and ::max() moduleStart[i] = ::max(0U, moduleStart[i]); @@ -535,7 +535,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(gpuClusteringConstants::maxNumModules, stream); nModules_Clusters_h = cms::cuda::make_host_unique(2, stream); @@ -585,10 +585,10 @@ namespace pixelgpudetails { { // clusterizer ... - using namespace gpuClustering; + using namespace gpuClusteringConstants; int threadsPerBlock = 256; int blocks = - (std::max(int(wordCounter), int(gpuClustering::maxNumModules)) + threadsPerBlock - 1) / threadsPerBlock; + (std::max(int(wordCounter), int(gpuClusteringConstants::maxNumModules)) + threadsPerBlock - 1) / threadsPerBlock; gpuCalibPixel::calibDigis<<>>(isRun2, digis_d.moduleInd(), @@ -658,7 +658,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() + gpuClusteringConstants::maxNumModules, sizeof(uint32_t), cudaMemcpyDefault, stream)); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h index c21c792f39c30..d5eb5cc71d63e 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h @@ -10,7 +10,7 @@ namespace gpuCalibPixel { - using gpuClustering::invalidModuleId; + using gpuClusteringConstants::invalidModuleId; // valid for run2 constexpr float VCaltoElectronGain = 47; // L2-4: 47 +- 4.7 @@ -34,7 +34,7 @@ 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 < gpuClusteringConstants::maxNumModules; i += gridDim.x * blockDim.x) { nClustersInModule[i] = 0; } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h index d9520da80b695..9bd64816f01e6 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -8,7 +8,7 @@ #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "HeterogeneousCore/CUDAUtilities/interface/prefixScan.h" -namespace gpuClustering { +namespace gpuClusteringConstants { __global__ void clusterChargeCut( uint16_t* __restrict__ id, // module id of each pixel (modified if bad cluster) @@ -120,6 +120,6 @@ namespace gpuClustering { } // loop on modules } -} // namespace gpuClustering +} // namespace gpuClusteringConstants #endif // RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterChargeCut_h diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h index 8467c1d74e2d9..b159aec149eb6 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h @@ -9,7 +9,7 @@ #include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" -namespace gpuClustering { +namespace gpuClusteringConstants { #ifdef GPU_DEBUG __device__ uint32_t gMaxHit = 0; @@ -298,6 +298,6 @@ namespace gpuClustering { } } // module loop } -} // namespace gpuClustering +} // namespace gpuClusteringConstants #endif // RecoLocalTracker_SiPixelClusterizer_plugins_gpuClustering_h diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h index 02611ab1cac1d..bc9236692ce51 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h +++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h @@ -24,7 +24,7 @@ int main(void) { cms::cudatest::requireDevices(); #endif // __CUDACC__ - using namespace gpuClustering; + using namespace gpuClusteringConstants; constexpr int numElements = 256 * maxNumModules; diff --git a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h index 23f9056cc1d67..ed447f8835997 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h @@ -98,7 +98,7 @@ namespace pixelCPEforGPU { int16_t ysize[N]; }; - constexpr int32_t MaxHitsInIter = gpuClustering::maxHitsInIter; + constexpr int32_t MaxHitsInIter = gpuClusteringConstants::maxHitsInIter; using ClusParams = ClusParamsT; constexpr inline void computeAnglesFromDet( diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc index 8c16be54e5774..06b0a88cb5108 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitConverter.cc @@ -191,18 +191,18 @@ namespace cms { const edmNew::DetSetVector& input = *inputhandle; // allocate a buffer for the indices of the clusters - auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); + auto hmsp = std::make_unique(gpuClusteringConstants::maxNumModules + 1); // hitsModuleStart is a non-owning pointer to the buffer auto hitsModuleStart = hmsp.get(); // fill cluster arrays - std::array clusInModule{}; + std::array clusInModule{}; 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)) + if (gind >= int(gpuClusteringConstants::maxNumModules)) continue; auto const nclus = dsv.size(); assert(nclus > 0); @@ -210,10 +210,10 @@ namespace cms { numberOfClusters += nclus; } hitsModuleStart[0] = 0; - assert(clusInModule.size() > gpuClustering::maxNumModules); + assert(clusInModule.size() > gpuClusteringConstants::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[gpuClusteringConstants::maxNumModules])); // 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)); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc index 6b717175c31bc..c38e52b40fb39 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromCUDA.cc @@ -84,8 +84,8 @@ void SiPixelRecHitFromCUDA::acquire(edm::Event const& iEvent, 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()); + auto hmsp = std::make_unique(gpuClusteringConstants::maxNumModules + 1); + std::copy(hitsModuleStart_.get(), hitsModuleStart_.get() + gpuClusteringConstants::maxNumModules + 1, hmsp.get()); // 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)); @@ -105,7 +105,7 @@ void SiPixelRecHitFromCUDA::produce(edm::Event& iEvent, edm::EventSetup const& e edm::Handle hclusters = iEvent.getHandle(clusterToken_); auto const& input = *hclusters; - constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule; + constexpr uint32_t maxHitsInModule = gpuClusteringConstants::maxHitsInModule; int numberOfDetUnits = 0; int numberOfClusters = 0; diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc index b9bdb9a5fa7a1..2133562cef1aa 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc @@ -33,7 +33,7 @@ class SiPixelRecHitSoAFromLegacy : public edm::global::EDProducer<> { static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); - using HitModuleStart = std::array; + using HitModuleStart = std::array; using HMSstorage = HostProduct; private: @@ -90,7 +90,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv auto const& input = *hclusters; // allocate a buffer for the indices of the clusters - auto hmsp = std::make_unique(gpuClustering::maxNumModules + 1); + auto hmsp = std::make_unique(gpuClusteringConstants::maxNumModules + 1); // hitsModuleStart is a non-owning pointer to the buffer auto hitsModuleStart = hmsp.get(); // wrap the buffer in a HostProduct @@ -110,13 +110,13 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv std::vector, SiPixelCluster>> clusterRef; - constexpr uint32_t maxHitsInModule = gpuClustering::maxHitsInModule; + constexpr uint32_t maxHitsInModule = gpuClusteringConstants::maxHitsInModule; HitModuleStart moduleStart_; // index of the first pixel of each module HitModuleStart clusInModule_; memset(&clusInModule_, 0, sizeof(HitModuleStart)); // needed?? - assert(gpuClustering::maxNumModules + 1 == clusInModule_.size()); - assert(0 == clusInModule_[gpuClustering::maxNumModules]); + assert(gpuClusteringConstants::maxNumModules + 1 == clusInModule_.size()); + assert(0 == clusInModule_[gpuClusteringConstants::maxNumModules]); uint32_t moduleId_; moduleStart_[1] = 0; // we run sequentially.... @@ -130,7 +130,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 < gpuClustering::maxNumModules); + assert(gind < gpuClusteringConstants::maxNumModules); auto const nclus = dsv.size(); clusInModule_[gind] = nclus; numberOfClusters += nclus; @@ -138,7 +138,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[gpuClustering::maxNumModules])); + assert(numberOfClusters == int(hitsModuleStart[gpuClusteringConstants::maxNumModules])); // output SoA auto output = std::make_unique(numberOfClusters, &cpeView, hitsModuleStart, nullptr); @@ -151,7 +151,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv } if (convert2Legacy_) - legacyOutput->reserve(gpuClustering::maxNumModules, numberOfClusters); + legacyOutput->reserve(gpuClusteringConstants::maxNumModules, numberOfClusters); int numberOfDetUnits = 0; int numberOfHits = 0; @@ -161,7 +161,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 < gpuClustering::maxNumModules); + assert(gind < gpuClusteringConstants::maxNumModules); const PixelGeomDetUnit* pixDet = dynamic_cast(genericDet); assert(pixDet); auto const nclus = dsv.size(); @@ -217,7 +217,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv if (h - fc < maxHitsInModule) assert(gind == output->view()->detectorIndex(h)); else - assert(gpuClustering::invalidModuleId == output->view()->detectorIndex(h)); + assert(gpuClusteringConstants::invalidModuleId == output->view()->detectorIndex(h)); if (convert2Legacy_) { SiPixelRecHitCollectionNew::FastFiller recHitsOnDetUnit(*legacyOutput, detid); for (auto h = fc; h < lc; ++h) { diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index 2401fed6c5171..f0255aaa21aba 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -54,7 +54,7 @@ namespace gpuPixelRecHits { } // to be moved in common namespace... - using gpuClustering::invalidModuleId; + using gpuClusteringConstants::invalidModuleId; constexpr int32_t MaxHitsInIter = pixelCPEforGPU::MaxHitsInIter; using ClusParams = pixelCPEforGPU::ClusParams; diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackProducerFromSoA.cc index cdea22c3a8a24..4ec817582bb73 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/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h index 5c31c31ca62f6..d0afc9bc15a59 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h @@ -57,7 +57,7 @@ namespace CAConstants { using OuterHitOfCell = cms::cuda::VecArray; using TuplesContainer = cms::cuda::OneToManyAssoc; using HitToTuple = - cms::cuda::OneToManyAssoc; // 3.5 should be enough + cms::cuda::OneToManyAssoc; // 3.5 should be enough using TupleMultiplicity = cms::cuda::OneToManyAssoc; } // namespace CAConstants diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc index a4f032d21cc41..46c5f57aa4f2a 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc @@ -78,7 +78,7 @@ void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA * cms::cuda::launchZero(tuples_d, cudaStream); auto nhits = hh.nHits(); - assert(nhits <= gpuClustering::maxNumberOfHits); + assert(nhits <= gpuClusteringConstants::maxNumberOfHits); // std::cout << "N hits " << nhits << std::endl; // if (nhits<2) std::cout << "too few hits " << nhits << std::endl; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu index edba7d9cb746a..1ad717da49d89 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu @@ -24,7 +24,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA * cms::cuda::launchZero(tuples_d, cudaStream); auto nhits = hh.nHits(); - assert(nhits <= gpuClustering::maxNumberOfHits); + assert(nhits <= gpuClusteringConstants::maxNumberOfHits); // std::cout << "N hits " << nhits << std::endl; // if (nhits<2) std::cout << "too few hits " << nhits << std::endl; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h b/RecoPixelVertexing/PixelTriplets/plugins/gpuPixelDoubletsAlgos.h index 5c0d5a252b684..e2a770168052d 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 > gpuClustering::maxNumModules) + if (mi > gpuClusteringConstants::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 > gpuClustering::maxNumModules) + if (mo > gpuClusteringConstants::maxNumModules) continue; // invalid if (doZ0Cut && z0cutoff(oi)) diff --git a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu index 0aab26d9cc091..b70e4296a0b7b 100644 --- a/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu +++ b/SimTracker/TrackerHitAssociation/plugins/ClusterSLOnGPU.cu @@ -21,8 +21,8 @@ __global__ void simLink(const SiPixelDigisCUDA::DeviceConstView* dd, ClusterSLView sl, uint32_t n) { constexpr uint32_t invTK = 0; // std::numeric_limits::max(); - using gpuClustering::invalidModuleId; - using gpuClustering::maxNumModules; + using gpuClusteringConstants::invalidModuleId; + using gpuClusteringConstants::maxNumModules; auto const& hh = *hhp; auto i = blockIdx.x * blockDim.x + threadIdx.x; From 98e0ec27c57b2e73f222a97fcc8d75d0a4a6ed38 Mon Sep 17 00:00:00 2001 From: Angela Czirkos Date: Fri, 12 Mar 2021 10:35:54 +0100 Subject: [PATCH 5/5] Fix RecoLocalTracker/SiPixelClusterizer build file --- RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml | 1 - 1 file changed, 1 deletion(-) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml b/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml index 8c9e1456eb1bb..82519c629f620 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/BuildFile.xml @@ -9,7 +9,6 @@ -