Merge pull request #35317 from czangela/gpu_crash_34831_fix_2021_09_1…

…0_backport_12_0_X [Backport 12_0_X] Fix for SiPixelRecHitFromCUDA crash during online GPU tests
cms-sw · Sep 20, 2021 · db48cba · db48cba
2 parents b9696bc + fc76746
commit db48cba
Show file tree

Hide file tree

Showing 4 changed files with 15 additions and 4 deletions.
diff --git a/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h b/CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h
@@ -18,6 +18,7 @@ namespace gpuClustering {
   constexpr uint16_t maxNumModules = 2000;
   constexpr int32_t maxNumClustersPerModules = maxHitsInModule();
   constexpr uint16_t invalidModuleId = std::numeric_limits<uint16_t>::max() - 1;
+  constexpr int invalidClusterId = -9999;
   static_assert(invalidModuleId > maxNumModules);  // invalidModuleId must be > maxNumModules
 
 }  // namespace gpuClustering

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc
@@ -81,8 +81,14 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con
   edm::DetSet<PixelDigi>* detDigis = nullptr;
   uint32_t detId = 0;
   for (uint32_t i = 0; i < nDigis; i++) {
-    if (digis.pdigi(i) == 0)
+    // check for uninitialized digis
+    // this is set in RawToDigi_kernel in SiPixelRawToClusterGPUKernel.cu
+    if (digis.rawIdArr(i) == 0)
       continue;
+    // check for noisy/dead pixels (electrons set to 0)
+    if (digis.adc(i) == 0)
+      continue;
+
     detId = digis.rawIdArr(i);
     if (storeDigis_) {
       detDigis = &collection->find_or_insert(detId);
@@ -134,7 +140,11 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con
   };
 
   for (uint32_t i = 0; i < nDigis; i++) {
-    if (digis.pdigi(i) == 0)
+    // check for uninitialized digis
+    if (digis.rawIdArr(i) == 0)
+      continue;
+    // check for noisy/dead pixels (electrons set to 0)
+    if (digis.adc(i) == 0)
       continue;
     if (digis.clus(i) > 9000)
       continue;  // not in cluster; TODO add an assert for the size

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h
@@ -54,9 +54,9 @@ namespace gpuCalibPixel {
       float gain = ret.second;
       // float pedestal = 0; float gain = 1.;
       if (isDeadColumn | isNoisyColumn) {
+        printf("bad pixel at %d in %d\n", i, id[i]);
         id[i] = invalidModuleId;
         adc[i] = 0;
-        printf("bad pixel at %d in %d\n", i, id[i]);
       } else {
         float vcal = adc[i] * gain - pedestal * gain;
         adc[i] = std::max(100, int(vcal * conversionFactor + offset));

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h
@@ -275,7 +275,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] == invalidModuleId) {  // skip invalid pixels
-          clusterId[i] = -9999;
+          clusterId[i] = invalidClusterId;
           continue;
         }
         clusterId[i] = -clusterId[i] - 1;