Address more HCAL review comments (#568)

Further simplify HCAL raw data template specialisations. Clean up commented out code.
cms-patatrack · Nov 12, 2020 · a6bb180 · a6bb180
1 parent 2b9d2a8
commit a6bb180
Show file tree

Hide file tree

Showing 7 changed files with 62 additions and 140 deletions.
diff --git a/CUDADataFormats/HcalDigi/interface/DigiCollection.h b/CUDADataFormats/HcalDigi/interface/DigiCollection.h
@@ -7,44 +7,34 @@ namespace hcal {
 
   // FLAVOR_HE_QIE11 = 1; Phase1 upgrade
   struct Flavor1 {
-    using adc_type = uint8_t;
-    using tdc_type = uint8_t;
-    using soibit_type = uint8_t;
-
     static constexpr int WORDS_PER_SAMPLE = 1;
     static constexpr int SAMPLES_PER_WORD = 1;
     static constexpr int HEADER_WORDS = 1;
 
-    static constexpr adc_type adc(uint16_t const* const sample_start) { return (*sample_start & 0xff); }
-    static constexpr tdc_type tdc(uint16_t const* const sample_start) { return (*sample_start >> 8) & 0x3f; }
-    static constexpr soibit_type soibit(uint16_t const* const sample_start) { return (*sample_start >> 14) & 0x1; }
+    static constexpr uint8_t adc(uint16_t const* const sample_start) { return (*sample_start & 0xff); }
+    static constexpr uint8_t tdc(uint16_t const* const sample_start) { return (*sample_start >> 8) & 0x3f; }
+    static constexpr uint8_t soibit(uint16_t const* const sample_start) { return (*sample_start >> 14) & 0x1; }
   };
 
   // FLAVOR_HB_QIE11 = 3; Phase1 upgrade
   struct Flavor3 {
-    using adc_type = uint8_t;
-    using tdc_type = uint8_t;
-    using soibit_type = uint8_t;
-
     static constexpr int WORDS_PER_SAMPLE = 1;
     static constexpr int SAMPLES_PER_WORD = 1;
     static constexpr int HEADER_WORDS = 1;
 
-    static constexpr adc_type adc(uint16_t const* const sample_start) { return (*sample_start & 0xff); }
-    static constexpr tdc_type tdc(uint16_t const* const sample_start) { return ((*sample_start >> 8) & 0x3); }
-    static constexpr soibit_type soibit(uint16_t const* const sample_start) { return ((*sample_start >> 14) & 0x1); }
+    static constexpr uint8_t adc(uint16_t const* const sample_start) { return (*sample_start & 0xff); }
+    static constexpr uint8_t tdc(uint16_t const* const sample_start) { return ((*sample_start >> 8) & 0x3); }
+    static constexpr uint8_t soibit(uint16_t const* const sample_start) { return ((*sample_start >> 14) & 0x1); }
     static constexpr uint8_t capid(uint16_t const* const sample_start) { return ((*sample_start >> 10) & 0x3); }
   };
 
   // FLAVOR_HB_QIE10 = 5; Phase0
   struct Flavor5 {
-    using adc_type = uint8_t;
-
     static constexpr float WORDS_PER_SAMPLE = 0.5;
     static constexpr int SAMPLES_PER_WORD = 2;
     static constexpr int HEADER_WORDS = 1;
 
-    static constexpr adc_type adc(uint16_t const* const sample_start, uint8_t const shifter) {
+    static constexpr uint8_t adc(uint16_t const* const sample_start, uint8_t const shifter) {
       return ((*sample_start >> shifter * 8) & 0x7f);
     }
   };
@@ -61,22 +51,22 @@ namespace hcal {
   }
 
   template <typename Flavor>
-  constexpr typename Flavor::soibit_type soibit_for_sample(uint16_t const* const dfstart, uint32_t const sample) {
+  constexpr uint8_t soibit_for_sample(uint16_t const* const dfstart, uint32_t const sample) {
     return Flavor::soibit(dfstart + Flavor::HEADER_WORDS + sample * Flavor::WORDS_PER_SAMPLE);
   }
 
   template <typename Flavor>
-  constexpr typename Flavor::adc_type adc_for_sample(uint16_t const* const dfstart, uint32_t const sample) {
+  constexpr uint8_t adc_for_sample(uint16_t const* const dfstart, uint32_t const sample) {
     return Flavor::adc(dfstart + Flavor::HEADER_WORDS + sample * Flavor::WORDS_PER_SAMPLE);
   }
 
   template <typename Flavor>
-  constexpr typename Flavor::tdc_type tdc_for_sample(uint16_t const* const dfstart, uint32_t const sample) {
+  constexpr uint8_t tdc_for_sample(uint16_t const* const dfstart, uint32_t const sample) {
     return Flavor::tdc(dfstart + Flavor::HEADER_WORDS + sample * Flavor::WORDS_PER_SAMPLE);
   }
 
   template <>
-  constexpr Flavor5::adc_type adc_for_sample<Flavor5>(uint16_t const* const dfstart, uint32_t const sample) {
+  constexpr uint8_t adc_for_sample<Flavor5>(uint16_t const* const dfstart, uint32_t const sample) {
     // avoid using WORDS_PER_SAMPLE and simply shift
     return Flavor5::adc(dfstart + Flavor5::HEADER_WORDS + (sample >> 1), sample % 2);
   }
@@ -88,12 +78,10 @@ namespace hcal {
 
   template <typename Flavor>
   constexpr uint32_t compute_nsamples(uint32_t const nwords) {
-    return (nwords - Flavor::HEADER_WORDS) / Flavor::WORDS_PER_SAMPLE;
-  }
-
-  template <>
-  constexpr uint32_t compute_nsamples<Flavor5>(uint32_t const nwords) {
-    return (nwords - Flavor5::HEADER_WORDS) * Flavor5::SAMPLES_PER_WORD;
+    if constexpr (Flavor::SAMPLES_PER_WORD >= 1)
+      return (nwords - Flavor::HEADER_WORDS) * Flavor::SAMPLES_PER_WORD;
+    else
+      return (nwords - Flavor::HEADER_WORDS) / Flavor::WORDS_PER_SAMPLE;
   }
 
   //
@@ -138,12 +126,7 @@ namespace hcal {
   template <typename StoragePolicy>
   struct DigiCollection<Flavor5, StoragePolicy> : public DigiCollectionBase<StoragePolicy> {
     DigiCollection() = default;
-    //DigiCollection(
-    //        uint32_t *ids, uint16_t *data, uint8_t *presamples,
-    //        uint32_t ndigis, uint32_t stride)
-    //    : DigiCollectionBase(ids, data, ndigis, stride)
-    //    , npresamples{npresamples}
-    //{}
+
     DigiCollection(DigiCollection const&) = default;
     DigiCollection& operator=(DigiCollection const&) = default;
 

diff --git a/EventFilter/HcalRawToDigi/bin/makeHcalRaw2DigiGpuValidationPlots.cpp b/EventFilter/HcalRawToDigi/bin/makeHcalRaw2DigiGpuValidationPlots.cpp
@@ -7,12 +7,12 @@
 #include <TFile.h>
 #include <TH1D.h>
 #include <TH2D.h>
-#include <TTree.h>
 #include <TPaveStats.h>
+#include <TTree.h>
 
+#include "CUDADataFormats/HcalDigi/interface/DigiCollection.h"
 #include "DataFormats/Common/interface/Wrapper.h"
 #include "DataFormats/HcalDigi/interface/HcalDigiCollections.h"
-#include "CUDADataFormats/HcalDigi/interface/DigiCollection.h"
 
 #define CREATE_HIST_1D(varname, nbins, first, last) auto varname = new TH1D(#varname, #varname, nbins, first, last)
 
@@ -103,27 +103,6 @@ int main(int argc, char* argv[]) {
   CREATE_HIST_2D(hTDCf01HEGPUvsCPU, 64, 0, 64);
   CREATE_HIST_2D(hTDCf3HBGPUvsCPU, 4, 0, 4);
 
-  /*
-    auto hADCEBGPU = new TH1D("hADCEBGPU", "hADCEBGPU", nbins, 0, last);
-    auto hADCEBCPU = new TH1D("hADCEBCPU", "hADCEBCPU", nbins, 0, last);
-    auto hADCEEGPU = new TH1D("hADCEEGPU", "hADCEEGPU", nbins, 0, last);
-    auto hADCEECPU = new TH1D("hADCEECPU", "hADCEECPU", nbins, 0, last);
-
-    auto hGainEBGPU = new TH1D("hGainEBGPU", "hGainEBGPU", 4, 0, 4);
-    auto hGainEBCPU = new TH1D("hGainEBCPU", "hGainEBCPU", 4, 0, 4);
-    auto hGainEEGPU = new TH1D("hGainEEGPU", "hGainEEGPU", 4, 0, 4);
-    auto hGainEECPU = new TH1D("hGainEECPU", "hGainEECPU", 4, 0, 4);
-
-    auto hADCEBGPUvsCPU = new TH2D("hADCEBGPUvsCPU", "hADCEBGPUvsCPU",
-        nbins, 0, last, nbins, 0, last);
-    auto hADCEEGPUvsCPU = new TH2D("hADCEEGPUvsCPU", "hADCEEGPUvsCPU",
-        nbins, 0, last, nbins, 0, last);
-    auto hGainEBGPUvsCPU = new TH2D("hGainEBGPUvsCPU", "hGainEBGPUvsCPU",
-        4, 0, 4, 4, 0, 4);
-    auto hGainEEGPUvsCPU = new TH2D("hGainEEGPUvsCPU", "hGainEEGPUvsCPU",
-        4, 0, 4, 4, 0, 4);
-        */
-
   // prep input
   TFile rfin{inFileName.c_str()};
   TTree* rt = (TTree*)rfin.Get("Events");
@@ -313,7 +292,7 @@ int main(int argc, char* argv[]) {
           hADCf5HBGPUvsCPU->Fill(cpuadc, gpuadc);
 
           // the must for us at RAW Decoding stage
-          assert(static_cast<hcal::Flavor5::adc_type>(cpuadc) == gpuadc);
+          assert(static_cast<uint8_t>(cpuadc) == gpuadc);
           assert(static_cast<uint8_t>(cpucapid) == gpucapid);
         }
       }

diff --git a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
@@ -5,7 +5,6 @@
 
 #include "CUDADataFormats/HcalDigi/interface/DigiCollection.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
 
@@ -45,8 +44,8 @@ namespace hcal {
     };
 
     struct ScratchDataGPU {
-      // depends on tHE number of output collections
-      // that is a statically known predefined number!!!
+      // depends on the number of output collections
+      // that is a statically known predefined number
       cms::cuda::device::unique_ptr<uint32_t[]> pChannelsCounters;
     };
 
@@ -58,27 +57,16 @@ namespace hcal {
       void allocate(ConfigurationParameters const &config, cudaStream_t cudaStream) {
         digisF01HE.data = cms::cuda::make_device_unique<uint16_t[]>(
             config.maxChannelsF01HE * compute_stride<Flavor1>(config.nsamplesF01HE), cudaStream);
-        //cudaCheck(
-        //    cudaMalloc((void **)&digisF01HE.data,
-        //               config.maxChannelsF01HE * sizeof(uint16_t) * compute_stride<Flavor1>(config.nsamplesF01HE)));
         digisF01HE.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannelsF01HE, cudaStream);
-        //cudaCheck(cudaMalloc((void **)&digisF01HE.ids, sizeof(uint32_t) * config.maxChannelsF01HE));
 
         digisF5HB.data = cms::cuda::make_device_unique<uint16_t[]>(
             config.maxChannelsF5HB * compute_stride<Flavor5>(config.nsamplesF5HB), cudaStream);
-        //cudaCheck(cudaMalloc((void **)&digisF5HB.data,
-        //                     config.maxChannelsF5HB * sizeof(uint16_t) * compute_stride<Flavor5>(config.nsamplesF5HB)));
         digisF5HB.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannelsF5HB, cudaStream);
-        //cudaCheck(cudaMalloc((void **)&digisF5HB.ids, sizeof(uint32_t) * config.maxChannelsF5HB));
         digisF5HB.npresamples = cms::cuda::make_device_unique<uint8_t[]>(config.maxChannelsF5HB, cudaStream);
-        //cudaCheck(cudaMalloc((void **)&digisF5HB.npresamples, sizeof(uint8_t) * config.maxChannelsF5HB));
 
         digisF3HB.data = cms::cuda::make_device_unique<uint16_t[]>(
             config.maxChannelsF3HB * compute_stride<Flavor3>(config.nsamplesF3HB), cudaStream);
-        //cudaCheck(cudaMalloc((void **)&digisF3HB.data,
-        //                     config.maxChannelsF3HB * sizeof(uint16_t) * compute_stride<Flavor3>(config.nsamplesF3HB)));
         digisF3HB.ids = cms::cuda::make_device_unique<uint32_t[]>(config.maxChannelsF3HB, cudaStream);
-        //cudaCheck(cudaMalloc((void **)&digisF3HB.ids, config.maxChannelsF3HB * sizeof(uint32_t)));
       }
     };
 

diff --git a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
@@ -164,12 +164,10 @@ namespace hcal {
 
       // get to the payload
       auto const* payload64 = buffer + 2 + namc + amcoffset;
-      //amcoffset += amcSize;
 
 #ifdef HCAL_RAWDECODE_GPUDEBUG
       // uhtr header v1 1st 64 bits
       auto const payload64_w0 = payload64[0];
-      //uint32_t const data_length64 = payload64_w0 & 0xfffff;
 #endif
       // uhtr n bytes comes from amcSize, according to the cpu version!
       uint32_t const data_length64 = amcSize;
@@ -211,12 +209,11 @@ namespace hcal {
       // skip uhtr header words
       auto const channelDataSize = data_length64 - 2;        // 2 uhtr header v1 words
       auto const* channelDataBuffer64Start = payload64 + 2;  // 2 uhtr header v2 wds
-      //auto const* channelDataBuffer64End = channelDataBuffer64Start + channelDataSize;
       auto const* ptr = reinterpret_cast<uint16_t const*>(channelDataBuffer64Start);
       auto const* end = ptr + sizeof(uint64_t) / sizeof(uint16_t) * (channelDataSize - 1);
       auto const t_rank = thread_group.thread_rank();
 
-      // iterate thru the channel data
+      // iterate through the channel data
       while (ptr != end) {
         // this is the starting point for this thread group for this iteration
         // with respect to this pointer every thread will move forward afterwards
@@ -237,8 +234,7 @@ namespace hcal {
           if (is_channel_header_word(ptr))
             ++ptr;
           else {
-            // go to the next channel and do not consider this guy as a
-            // channel
+            // go to the next channel and do not consider this guy as a channel
             while (ptr != end)
               if (!is_channel_header_word(ptr))
                 ++ptr;
@@ -261,18 +257,7 @@ namespace hcal {
         printf("ptr - start_ptr = %d counter = %d rank = %d\n", static_cast<int>(ptr - start_ptr), counter, t_rank);
 #endif
 
-        // assume that if all is valid, ptr points
-        // to the header word of the channel to be decoded
-        // skip to the next channel header word if above assumption
-        // does not hold
-        //uint8_t const fw_lastbit = (*ptr >> 15) & 0x1;
-        //if (fw_lastbit != 1) {
-        //    ptr++;
-        //    continue;
-        //}
-
-        // when the end is near, channels will land outside of the [start_ptr, end)
-        // region
+        // when the end is near, channels will land outside of the [start_ptr, end) region
         if (ptr != end) {
           // for all of the flavors, these 2 guys have the same bit layout
           uint8_t const flavor = (ptr[0] >> 12) & 0x7;
@@ -469,9 +454,6 @@ namespace hcal {
               uint32_t const nwords = channel_end - channel_header_word;
 
               // filter out this digi if nwords does not equal expected
-              //uint32_t const expected_words =
-              //    nsamplesF5HB * Flavor5::WORDS_PER_SAMPLE +
-              //    Flavor5::HEADER_WORDS;
               auto const expected_words = compute_stride<Flavor5>(nsamplesF5HB);
               if (nwords != expected_words)
                 break;
@@ -497,16 +479,14 @@ namespace hcal {
               HcalElectronicsId eid{uhtrcrate, uhtrslot, fiber, fchannel, false};
               auto const did = DetId{eid2did[eid.linearIndex()]};
 
-              /*
-                if (eid.rawId() >= HcalElectronicsId::maxLinearIndex) {
+              /* uncomment to check the linear index validity
+              if (eid.rawId() >= HcalElectronicsId::maxLinearIndex) {
 #ifdef HCAL_RAWDECODE_GPUDEBUG
-                    printf("*** rawid = %u has no known det id***\n",
-                        eid.rawId());
+                  printf("*** rawid = %u has no known det id***\n", eid.rawId());
 #endif
-                    break;
-                }
-                */
-              //auto const did = DetId{eid2did[eid.rawId()]};
+                  break;
+              }
+              */
 
 #ifdef HCAL_RAWDECODE_GPUDEBUG
               printf("erawId = %u linearIndex = %u drawid = %u\n", eid.rawId(), eid.linearIndex(), did.rawId());

diff --git a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
@@ -76,10 +76,6 @@ namespace hcal {
       float timeSigmaSiPM, timeSigmaHPD;
       float ts4Thresh;
 
-      std::vector<int> pulseOffsets;
-      // FIXME remove pulseOffsets - they come from esproduce now
-      //int* pulseOffsetsDevice = nullptr;
-
       std::array<uint32_t, 3> kernelMinimizeThreads;
 
       // FIXME:
@@ -113,7 +109,7 @@ namespace hcal {
       DigiCollection<Flavor3, ::calo::common::DevStoragePolicy> const& f3HBDigis;
     };
 
-  }  // namespace mahi
-}  // namespace reconstruction
+  }  // namespace reconstruction
+}  // namespace hcal
 
 #endif  // RecoLocalCalo_HcalRecProducers_src_DeclsForKernels_h
diff --git a/RecoLocalCalo/HcalRecProducers/src/HcalESProducersGPUDefs.cc b/RecoLocalCalo/HcalRecProducers/src/HcalESProducersGPUDefs.cc
@@ -1,44 +1,42 @@
-#include "HcalESProducerGPU.h"
-
-#include "CondFormats/HcalObjects/interface/HcalRecoParams.h"
-#include "CondFormats/HcalObjects/interface/HcalPedestals.h"
-#include "CondFormats/HcalObjects/interface/HcalGains.h"
-#include "CondFormats/HcalObjects/interface/HcalLUTCorrs.h"
-#include "CondFormats/HcalObjects/interface/HcalRespCorrs.h"
-#include "CondFormats/HcalObjects/interface/HcalTimeCorrs.h"
-#include "CondFormats/HcalObjects/interface/HcalPedestalWidths.h"
-#include "CondFormats/HcalObjects/interface/HcalGainWidths.h"
-#include "CondFormats/HcalObjects/interface/HcalQIEData.h"
-#include "CondFormats/HcalObjects/interface/HcalQIETypes.h"
-#include "CondFormats/HcalObjects/interface/HcalSiPMParameters.h"
-#include "CondFormats/HcalObjects/interface/HcalSiPMCharacteristics.h"
-
-#include "CondFormats/DataRecord/interface/HcalRecoParamsRcd.h"
-#include "CondFormats/DataRecord/interface/HcalPedestalsRcd.h"
+#include "CondFormats/DataRecord/interface/HcalGainWidthsRcd.h"
 #include "CondFormats/DataRecord/interface/HcalGainsRcd.h"
 #include "CondFormats/DataRecord/interface/HcalLUTCorrsRcd.h"
-#include "CondFormats/DataRecord/interface/HcalRespCorrsRcd.h"
-#include "CondFormats/DataRecord/interface/HcalTimeCorrsRcd.h"
 #include "CondFormats/DataRecord/interface/HcalPedestalWidthsRcd.h"
-#include "CondFormats/DataRecord/interface/HcalGainWidthsRcd.h"
+#include "CondFormats/DataRecord/interface/HcalPedestalsRcd.h"
 #include "CondFormats/DataRecord/interface/HcalQIEDataRcd.h"
 #include "CondFormats/DataRecord/interface/HcalQIETypesRcd.h"
-#include "CondFormats/DataRecord/interface/HcalSiPMParametersRcd.h"
+#include "CondFormats/DataRecord/interface/HcalRecoParamsRcd.h"
+#include "CondFormats/DataRecord/interface/HcalRespCorrsRcd.h"
 #include "CondFormats/DataRecord/interface/HcalSiPMCharacteristicsRcd.h"
-
-#include "CondFormats/HcalObjects/interface/HcalRecoParamsGPU.h"
-#include "RecoLocalCalo/HcalRecAlgos/interface/HcalRecoParamsWithPulseShapesGPU.h"
-#include "CondFormats/HcalObjects/interface/HcalPedestalsGPU.h"
+#include "CondFormats/DataRecord/interface/HcalSiPMParametersRcd.h"
+#include "CondFormats/DataRecord/interface/HcalTimeCorrsRcd.h"
+#include "CondFormats/HcalObjects/interface/HcalGainWidths.h"
+#include "CondFormats/HcalObjects/interface/HcalGainWidthsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalGains.h"
 #include "CondFormats/HcalObjects/interface/HcalGainsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalLUTCorrs.h"
 #include "CondFormats/HcalObjects/interface/HcalLUTCorrsGPU.h"
-#include "CondFormats/HcalObjects/interface/HcalRespCorrsGPU.h"
-#include "CondFormats/HcalObjects/interface/HcalTimeCorrsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalPedestalWidths.h"
 #include "CondFormats/HcalObjects/interface/HcalPedestalWidthsGPU.h"
-#include "CondFormats/HcalObjects/interface/HcalGainWidthsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalPedestals.h"
+#include "CondFormats/HcalObjects/interface/HcalPedestalsGPU.h"
 #include "CondFormats/HcalObjects/interface/HcalQIECodersGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalQIEData.h"
+#include "CondFormats/HcalObjects/interface/HcalQIETypes.h"
 #include "CondFormats/HcalObjects/interface/HcalQIETypesGPU.h"
-#include "CondFormats/HcalObjects/interface/HcalSiPMParametersGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalRecoParams.h"
+#include "CondFormats/HcalObjects/interface/HcalRecoParamsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalRespCorrs.h"
+#include "CondFormats/HcalObjects/interface/HcalRespCorrsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalSiPMCharacteristics.h"
 #include "CondFormats/HcalObjects/interface/HcalSiPMCharacteristicsGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalSiPMParameters.h"
+#include "CondFormats/HcalObjects/interface/HcalSiPMParametersGPU.h"
+#include "CondFormats/HcalObjects/interface/HcalTimeCorrs.h"
+#include "CondFormats/HcalObjects/interface/HcalTimeCorrsGPU.h"
+#include "RecoLocalCalo/HcalRecAlgos/interface/HcalRecoParamsWithPulseShapesGPU.h"
+
+#include "HcalESProducerGPU.h"
 
 using HcalRecoParamsGPUESProducer = HcalESProducerGPU<HcalRecoParamsRcd, HcalRecoParamsGPU, HcalRecoParams>;