Merge pull request #39740 from fwyzard/fix_HCAL_GPU_rechit_producer_124x

Skip invalid channels in the GPU HCAL RecHit producer [12.4.x]
cms-sw · Oct 19, 2022 · bc6cca8 · bc6cca8
2 parents fc7c9a9 + 76bd7f9
commit bc6cca8
Show file tree

Hide file tree

Showing 2 changed files with 20 additions and 4 deletions.
diff --git a/RecoLocalCalo/HcalRecProducers/src/HcalCPURecHitsProducer.cc b/RecoLocalCalo/HcalRecProducers/src/HcalCPURecHitsProducer.cc
@@ -103,14 +103,18 @@ void HcalCPURecHitsProducer::produce(edm::Event& event, edm::EventSetup const& s
     // did not set size with ctor as there is no setter for did
     recHitsLegacy->reserve(tmpRecHits_.did.size());
     for (uint32_t i = 0; i < tmpRecHits_.did.size(); i++) {
+      // skip bad channels
+      if (tmpRecHits_.chi2[i] < 0)
+        continue;
+
+      // build a legacy rechit with the computed detid and MAHI energy
       recHitsLegacy->emplace_back(HcalDetId{tmpRecHits_.did[i]},
                                   tmpRecHits_.energy[i],
                                   0  // timeRising
       );
-
-      // update newly pushed guy
-      (*recHitsLegacy)[i].setChiSquared(tmpRecHits_.chi2[i]);
-      (*recHitsLegacy)[i].setRawEnergy(tmpRecHits_.energyM0[i]);
+      // update the legacy rechit with the Chi2 and M0 values
+      recHitsLegacy->back().setChiSquared(tmpRecHits_.chi2[i]);
+      recHitsLegacy->back().setRawEnergy(tmpRecHits_.energyM0[i]);
     }
 
     // put the legacy collection

diff --git a/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu b/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
@@ -142,6 +142,7 @@ namespace hcal {
         method0Time[gch] = 0;
         outputEnergy[gch] = 0;
         outputChi2[gch] = 0;
+        soiSamples[gch] = -1;
       }
 
 #ifdef HCAL_MAHI_GPUDEBUG
@@ -265,6 +266,16 @@ namespace hcal {
       int32_t const soi = gch < nchannelsf01HE
                               ? soiSamples[gch]
                               : (gch < nchannelsf015 ? npresamplesf5HB[gch - nchannelsf01HE] : soiSamples[gch]);
+
+      bool badSOI = (soi < 0 or soi >= nsamplesForCompute);
+      if (badSOI and sampleWithinWindow == 0) {
+#ifdef GPU_DEBUG
+        printf("Found HBHE channel %d with invalid SOI %d\n", gch, soi);
+#endif
+        // mark the channel as bad
+        outputChi2[gch] = -9999.f;
+      }
+
       //int32_t const soi = gch >= nchannelsf01HE
       //    ? npresamplesf5HB[gch - nchannelsf01HE]
       //    : soiSamples[gch];
@@ -365,6 +376,7 @@ namespace hcal {
       __syncthreads();
 
       // NOTE: must take soi, as values for that thread are used...
+      // NOTE: does not run if soi is bad, because it does not match any sampleWithinWindow
       if (sampleWithinWindow == soi) {
         auto const method0_energy = shrMethod0EnergyAccum[lch];
         auto const val = shrMethod0EnergySamplePair[lch];