Consistently use gpuClustering::maxNumModules

CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DHeterogeneous.cc

@@ -13,7 +13,7 @@ cms::cuda::host::unique_ptr<float[]> TrackingRecHit2DCUDA::localCoordToHostAsync
 
 template <>
 cms::cuda::host::unique_ptr<uint32_t[]> TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cudaStream_t stream) const {
-  auto ret = cms::cuda::make_host_unique<uint32_t[]>(2001, stream);
-  cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, 4 * 2001, cudaMemcpyDefault, stream));
+  auto ret = cms::cuda::make_host_unique<uint32_t[]>(gpuClustering::maxNumModules + 1, stream);
+  cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (gpuClustering::maxNumModules + 1), cudaMemcpyDefault, stream));
   return ret;
 }

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<Range, int> rangeAndCols_[2000];
+  std::pair<Range, int> rangeAndCols_[gpuClustering::maxNumModules];
 
   float minPed_, maxPed_, minGain_, maxGain_;
   float pedPrecision_, gainPrecision_;

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<edm::DetSetVector<PixelDigi>>();
   auto outputClusters = std::make_unique<SiPixelClusterCollectionNew>();
-  outputClusters->reserve(2000, nDigis / 4);
+  outputClusters->reserve(gpuClustering::maxNumModules, nDigis / 4);
 
   edm::DetSet<PixelDigi>* detDigis = nullptr;
   for (uint32_t i = 0; i < nDigis; i++) {

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<uint16_t[]>(numElements);
   auto h_x = std::make_unique<uint16_t[]>(numElements);

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,15 +128,15 @@ 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;
   }
   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<TrackingRecHit2DCPU>(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<const PixelGeomDetUnit*>(genericDet);
     assert(pixDet);
     auto const nclus = DSViter->size();

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))

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<int32_t>::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}};