From 9ca193e2a0a509c08d7aa5ecd1a57d1c6df97cf0 Mon Sep 17 00:00:00 2001
From: Viktor Khristenko <vdkhristenko1991@gmail.com>
Date: Fri, 19 Jun 2020 12:46:01 +0200
Subject: [PATCH 1/6] make scratch use caching alloc

---
 EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h |  9 ++++++---
 EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu      |  6 +++---
 .../HcalRawToDigi/plugins/HcalRawToDigiGPU.cc       | 13 +++++++++----
 3 files changed, 18 insertions(+), 10 deletions(-)

diff --git a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
index 606053edb6801..08fcf6fcdefd9 100644
--- a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
+++ b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
@@ -6,6 +6,7 @@
 #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 "ElectronicsMappingGPU.h"
 
@@ -53,17 +54,19 @@ namespace hcal {
     struct ScratchDataGPU {
       // depends on tHE number of output collections
       // that is a statically known predefined number!!!
-      uint32_t *pChannelsCounters = nullptr;
+      cms::cuda::device::unique_ptr<uint32_t[]> pChannelsCounters;
 
+      /*
       void allocate(ConfigurationParameters const &) {
         cudaCheck(cudaMalloc((void **)&pChannelsCounters, sizeof(uint32_t) * numOutputCollections));
-      }
+      }*/
 
+      /*
       void deallocate(ConfigurationParameters const &) {
         if (pChannelsCounters) {
           cudaCheck(cudaFree(pChannelsCounters));
         }
-      }
+      }*/
     };
 
     struct OutputDataGPU {
diff --git a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
index 1589ec6cb1661..9bc5b9975f3e7 100644
--- a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
+++ b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
@@ -574,7 +574,7 @@ namespace hcal {
                                 nfedsWithData * sizeof(uint32_t),
                                 cudaMemcpyHostToDevice,
                                 cudaStream));
-      cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters, 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
+      cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters.get(), 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
       cudaCheck(cudaMemcpyAsync(
           inputGPU.feds, inputCPU.feds.data(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));
 
@@ -591,7 +591,7 @@ namespace hcal {
                                                                            outputGPU.digisF5HB.npresamples,
                                                                            outputGPU.digisF3HB.data,
                                                                            outputGPU.digisF3HB.ids,
-                                                                           scratchGPU.pChannelsCounters,
+                                                                           scratchGPU.pChannelsCounters.get(),
                                                                            config.nsamplesF01HE,
                                                                            config.nsamplesF5HB,
                                                                            config.nsamplesF3HB,
@@ -599,7 +599,7 @@ namespace hcal {
       cudaCheck(cudaGetLastError());
 
       cudaCheck(cudaMemcpyAsync(outputCPU.nchannels.data(),
-                                scratchGPU.pChannelsCounters,
+                                scratchGPU.pChannelsCounters.get(),
                                 sizeof(uint32_t) * numOutputCollections,
                                 cudaMemcpyDeviceToHost,
                                 cudaStream));
diff --git a/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc b/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
index 0e3a1a0d3b1e3..bc952fc936ba7 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
@@ -12,6 +12,7 @@
 #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 
 #include "DeclsForKernels.h"
 #include "DecodeGPU.h"
@@ -45,7 +46,6 @@ class HcalRawToDigiGPU : public edm::stream::EDProducer<edm::ExternalWork> {
   hcal::raw::InputDataCPU inputCPU_;
   hcal::raw::InputDataGPU inputGPU_;
   hcal::raw::OutputDataGPU outputGPU_;
-  hcal::raw::ScratchDataGPU scratchGPU_;
   hcal::raw::OutputDataCPU outputCPU_;
 };
 
@@ -93,7 +93,6 @@ HcalRawToDigiGPU::HcalRawToDigiGPU(const edm::ParameterSet& ps)
 
     inputGPU_.allocate();
     outputGPU_.allocate(config_);
-    scratchGPU_.allocate(config_);
   }
 }
 
@@ -103,7 +102,6 @@ HcalRawToDigiGPU::~HcalRawToDigiGPU() {
   if (cs and cs->enabled()) {
     inputGPU_.deallocate();
     outputGPU_.deallocate(config_);
-    scratchGPU_.deallocate(config_);
   }
 }
 
@@ -125,6 +123,13 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,
   edm::Handle<FEDRawDataCollection> rawDataHandle;
   event.getByToken(rawDataToken_, rawDataHandle);
 
+  // scratch
+  hcal::raw::ScratchDataGPU scratchGPU = {
+      cms::cuda::make_device_unique<uint32_t[]>(
+        hcal::raw::numOutputCollections,
+        ctx.stream())
+  };
+
   // iterate over feds
   // TODO: another idea
   //   - loop over all feds to unpack and enqueue cuda memcpy
@@ -159,7 +164,7 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,
   hcal::raw::entryPoint(inputCPU_,
                         inputGPU_,
                         outputGPU_,
-                        scratchGPU_,
+                        scratchGPU,
                         outputCPU_,
                         conditions,
                         config_,

From c7f93923ad372f06f1fe2b61be28a245a91904a7 Mon Sep 17 00:00:00 2001
From: Viktor Khristenko <vdkhristenko1991@gmail.com>
Date: Fri, 19 Jun 2020 13:25:42 +0200
Subject: [PATCH 2/6] use caching allocator for input cpu

---
 .../HcalRawToDigi/plugins/DeclsForKernels.h   | 20 +++++--------
 .../HcalRawToDigi/plugins/DecodeGPU.cu        |  6 ++--
 .../HcalRawToDigi/plugins/HcalRawToDigiGPU.cc | 29 +++++++++++++------
 3 files changed, 30 insertions(+), 25 deletions(-)

diff --git a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
index 08fcf6fcdefd9..591bdc40a48ed 100644
--- a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
+++ b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
@@ -7,6 +7,7 @@
 #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"
 
 #include "ElectronicsMappingGPU.h"
 
@@ -34,6 +35,11 @@ namespace hcal {
     };
 
     struct InputDataCPU {
+      cms::cuda::host::unique_ptr<unsigned char[]> data;
+      cms::cuda::host::unique_ptr<uint32_t[]> offsets;
+      cms::cuda::host::unique_ptr<int[]> feds;
+
+      /*
       std::vector<unsigned char, cms::cuda::HostAllocator<unsigned char>> data;
       std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t>> offsets;
       std::vector<int, cms::cuda::HostAllocator<int>> feds;
@@ -42,7 +48,7 @@ namespace hcal {
         data.resize(utca_nfeds_max * sizeof(unsigned char) * nbytes_per_fed_max);
         offsets.resize(utca_nfeds_max, 0);
         feds.resize(utca_nfeds_max, 0);
-      }
+      }*/
     };
 
     struct OutputDataCPU {
@@ -55,18 +61,6 @@ namespace hcal {
       // depends on tHE number of output collections
       // that is a statically known predefined number!!!
       cms::cuda::device::unique_ptr<uint32_t[]> pChannelsCounters;
-
-      /*
-      void allocate(ConfigurationParameters const &) {
-        cudaCheck(cudaMalloc((void **)&pChannelsCounters, sizeof(uint32_t) * numOutputCollections));
-      }*/
-
-      /*
-      void deallocate(ConfigurationParameters const &) {
-        if (pChannelsCounters) {
-          cudaCheck(cudaFree(pChannelsCounters));
-        }
-      }*/
     };
 
     struct OutputDataGPU {
diff --git a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
index 9bc5b9975f3e7..35cb128ba3d15 100644
--- a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
+++ b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
@@ -568,15 +568,15 @@ namespace hcal {
                     uint32_t const nbytesTotal) {
       // transfer
       cudaCheck(cudaMemcpyAsync(
-          inputGPU.data, inputCPU.data.data(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
+          inputGPU.data, inputCPU.data.get(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
       cudaCheck(cudaMemcpyAsync(inputGPU.offsets,
-                                inputCPU.offsets.data(),
+                                inputCPU.offsets.get(),
                                 nfedsWithData * sizeof(uint32_t),
                                 cudaMemcpyHostToDevice,
                                 cudaStream));
       cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters.get(), 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
       cudaCheck(cudaMemcpyAsync(
-          inputGPU.feds, inputCPU.feds.data(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));
+          inputGPU.feds, inputCPU.feds.get(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));
 
       // 12 is the max number of modules per crate
       kernel_rawdecode_test<32><<<nfedsWithData, 12 * 32, 0, cudaStream>>>(inputGPU.data,
diff --git a/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc b/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
index bc952fc936ba7..ce5e85659ebdc 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
@@ -43,7 +43,6 @@ class HcalRawToDigiGPU : public edm::stream::EDProducer<edm::ExternalWork> {
 
   hcal::raw::ConfigurationParameters config_;
   // FIXME move this to use raii
-  hcal::raw::InputDataCPU inputCPU_;
   hcal::raw::InputDataGPU inputGPU_;
   hcal::raw::OutputDataGPU outputGPU_;
   hcal::raw::OutputDataCPU outputCPU_;
@@ -88,7 +87,6 @@ HcalRawToDigiGPU::HcalRawToDigiGPU(const edm::ParameterSet& ps)
   // reserve memory and call CUDA API functions only if CUDA is available
   edm::Service<CUDAService> cs;
   if (cs and cs->enabled()) {
-    inputCPU_.allocate();
     outputCPU_.allocate();
 
     inputGPU_.allocate();
@@ -125,9 +123,22 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,
 
   // scratch
   hcal::raw::ScratchDataGPU scratchGPU = {
-      cms::cuda::make_device_unique<uint32_t[]>(
-        hcal::raw::numOutputCollections,
-        ctx.stream())
+    cms::cuda::make_device_unique<uint32_t[]>(
+      hcal::raw::numOutputCollections,
+      ctx.stream())
+  };
+
+  // input cpu data
+  hcal::raw::InputDataCPU inputCPU = {
+    cms::cuda::make_host_unique<unsigned char[]>(
+      hcal::raw::utca_nfeds_max * hcal::raw::nbytes_per_fed_max,
+      ctx.stream()),
+    cms::cuda::make_host_unique<uint32_t[]>(
+      hcal::raw::utca_nfeds_max,
+      ctx.stream()),
+    cms::cuda::make_host_unique<int[]>(
+      hcal::raw::utca_nfeds_max,
+      ctx.stream())
   };
 
   // iterate over feds
@@ -151,17 +162,17 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,
 #endif
 
     // copy raw data into plain buffer
-    std::memcpy(inputCPU_.data.data() + currentCummOffset, data.data(), nbytes);
+    std::memcpy(inputCPU.data.get() + currentCummOffset, data.data(), nbytes);
     // set the offset in bytes from the start
-    inputCPU_.offsets[counter] = currentCummOffset;
-    inputCPU_.feds[counter] = fed;
+    inputCPU.offsets[counter] = currentCummOffset;
+    inputCPU.feds[counter] = fed;
 
     // this is the current offset into the vector
     currentCummOffset += nbytes;
     ++counter;
   }
 
-  hcal::raw::entryPoint(inputCPU_,
+  hcal::raw::entryPoint(inputCPU,
                         inputGPU_,
                         outputGPU_,
                         scratchGPU,

From a84d73103d874cf0f72c7c243d8e186fe7f802c7 Mon Sep 17 00:00:00 2001
From: Viktor Khristenko <vdkhristenko1991@gmail.com>
Date: Thu, 2 Jul 2020 17:36:23 +0200
Subject: [PATCH 3/6] using allocators for EventHFilter/HcalRawToDigi

---
 CUDADataFormats/HcalCommon/interface/Common.h |  16 ++
 .../HcalCommon/src/classes_def.xml            |   2 +
 CUDADataFormats/HcalDigi/src/classes_def.xml  |  12 ++
 .../HcalRecHitSoA/src/classes_def.xml         |   3 +
 .../HcalRawToDigi/plugins/DeclsForKernels.h   | 128 ++++++----------
 .../HcalRawToDigi/plugins/DecodeGPU.cu        |  28 ++--
 .../plugins/HcalCPUDigisProducer.cc           |  89 ++---------
 .../plugins/HcalDigisProducerGPU.cc           | 144 +++++++-----------
 .../HcalRawToDigi/plugins/HcalRawToDigiGPU.cc |  61 ++++----
 9 files changed, 192 insertions(+), 291 deletions(-)

diff --git a/CUDADataFormats/HcalCommon/interface/Common.h b/CUDADataFormats/HcalCommon/interface/Common.h
index 1a2592889a6c7..53d49d06a89d1 100644
--- a/CUDADataFormats/HcalCommon/interface/Common.h
+++ b/CUDADataFormats/HcalCommon/interface/Common.h
@@ -4,6 +4,7 @@
 #include <vector>
 
 #include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 
 namespace hcal {
   namespace common {
@@ -13,6 +14,7 @@ namespace hcal {
 
       struct Vec {};
       struct Ptr {};
+      struct DevPtr {};
 
     }  // namespace tags
 
@@ -24,6 +26,11 @@ namespace hcal {
       uint32_t size;
     };
 
+    template<>
+    struct AddSize<tags::DevPtr> {
+        uint32_t size;
+    };
+
     struct ViewStoragePolicy {
       using TagType = tags::Ptr;
 
@@ -33,6 +40,15 @@ namespace hcal {
       };
     };
 
+    struct DevStoragePolicy {
+      using TagType = tags::DevPtr;
+
+      template<typename T>
+      struct StorageSelector {
+          using type = cms::cuda::device::unique_ptr<T[]>;
+      };
+    };
+
     template <template <typename> typename Allocator = std::allocator>
     struct VecStoragePolicy {
       using TagType = tags::Vec;
diff --git a/CUDADataFormats/HcalCommon/src/classes_def.xml b/CUDADataFormats/HcalCommon/src/classes_def.xml
index 173576a43950b..88778b86aa7e1 100644
--- a/CUDADataFormats/HcalCommon/src/classes_def.xml
+++ b/CUDADataFormats/HcalCommon/src/classes_def.xml
@@ -1,6 +1,8 @@
 <lcgdict>
     <enum name="hcal::common::tags::Ptr"/>
     <enum name="hcal::common::tags::Vec"/>
+    <enum name="hcal::common::tags::DevPtr"/>
     <class name="hcal::common::AddSize<hcal::common::tags::Ptr>"/>
     <class name="hcal::common::AddSize<hcal::common::tags::Vec>"/>
+    <class name="hcal::common::AddSize<hcal::common::tags::DevPtr>"/>
 </lcgdict>
diff --git a/CUDADataFormats/HcalDigi/src/classes_def.xml b/CUDADataFormats/HcalDigi/src/classes_def.xml
index 33e9b28a49b49..abe1769d9eeba 100644
--- a/CUDADataFormats/HcalDigi/src/classes_def.xml
+++ b/CUDADataFormats/HcalDigi/src/classes_def.xml
@@ -20,11 +20,23 @@
     <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::ViewStoragePolicy>>" persistent="false" />
     <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>" persistent="false" />
             
+    <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>" persistent="false" />
+    <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor2, hcal::common::DevStoragePolicy>>" persistent="false" />
+    <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>" persistent="false" />
+    <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::DevStoragePolicy>>" persistent="false" />
+    <class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>" persistent="false" />
+            
     <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::ViewStoragePolicy>>>" persistent="false" />
     <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor2, hcal::common::ViewStoragePolicy>>>" persistent="false" />
     <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::ViewStoragePolicy>>>" persistent="false" />
     <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::ViewStoragePolicy>>>" persistent="false" />
     <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>>" persistent="false" />
+                
+    <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor2, hcal::common::DevStoragePolicy>>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::DevStoragePolicy>>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>>" persistent="false" />
 
     <class name="edm::Wrapper<hcal::DigiCollection<hcal::Flavor01, hcal::common::VecStoragePolicy<std::allocator>>>" />
     <class name="edm::Wrapper<hcal::DigiCollection<hcal::Flavor2, hcal::common::VecStoragePolicy<std::allocator>>>" />
diff --git a/CUDADataFormats/HcalRecHitSoA/src/classes_def.xml b/CUDADataFormats/HcalRecHitSoA/src/classes_def.xml
index f9b6eb4ecc130..ce291190011c7 100644
--- a/CUDADataFormats/HcalRecHitSoA/src/classes_def.xml
+++ b/CUDADataFormats/HcalRecHitSoA/src/classes_def.xml
@@ -7,4 +7,7 @@
 
     <class name="cms::cuda::Product<hcal::RecHitCollection<hcal::common::ViewStoragePolicy>>" persistent="false" />
     <class name="edm::Wrapper<cms::cuda::Product<hcal::RecHitCollection<hcal::common::ViewStoragePolicy>>>" persistent="false" />
+                
+    <class name="cms::cuda::Product<hcal::RecHitCollection<hcal::common::DevStoragePolicy>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<hcal::RecHitCollection<hcal::common::DevStoragePolicy>>>" persistent="false" />
 </lcgdict>
diff --git a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
index 591bdc40a48ed..309e75b728e76 100644
--- a/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
+++ b/EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
@@ -38,23 +38,10 @@ namespace hcal {
       cms::cuda::host::unique_ptr<unsigned char[]> data;
       cms::cuda::host::unique_ptr<uint32_t[]> offsets;
       cms::cuda::host::unique_ptr<int[]> feds;
-
-      /*
-      std::vector<unsigned char, cms::cuda::HostAllocator<unsigned char>> data;
-      std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t>> offsets;
-      std::vector<int, cms::cuda::HostAllocator<int>> feds;
-
-      void allocate() {
-        data.resize(utca_nfeds_max * sizeof(unsigned char) * nbytes_per_fed_max);
-        offsets.resize(utca_nfeds_max, 0);
-        feds.resize(utca_nfeds_max, 0);
-      }*/
     };
 
     struct OutputDataCPU {
-      std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t>> nchannels;
-
-      void allocate() { nchannels.resize(numOutputCollections); }
+      cms::cuda::host::unique_ptr<uint32_t[]> nchannels;
     };
 
     struct ScratchDataGPU {
@@ -64,74 +51,59 @@ namespace hcal {
     };
 
     struct OutputDataGPU {
-      DigiCollection<Flavor01, common::ViewStoragePolicy> digisF01HE;
-      DigiCollection<Flavor5, common::ViewStoragePolicy> digisF5HB;
-      DigiCollection<Flavor3, common::ViewStoragePolicy> digisF3HB;
-
-      // qie 11 HE
-      /*
-    uint16_t *digisF01HE = nullptr;
-    uint32_t *idsF01HE = nullptr;
-
-    // qie 8 HB
-    uint16_t *digisF5HB = nullptr;
-    uint32_t *idsF5HB = nullptr;
-    uint8_t *npresamplesF5HB = nullptr
-    */
-
-      void allocate(ConfigurationParameters const &config) {
-        cudaCheck(
-            cudaMalloc((void **)&digisF01HE.data,
-                       config.maxChannelsF01HE * sizeof(uint16_t) * compute_stride<Flavor01>(config.nsamplesF01HE)));
-        cudaCheck(cudaMalloc((void **)&digisF01HE.ids, sizeof(uint32_t) * config.maxChannelsF01HE));
-
-        cudaCheck(cudaMalloc((void **)&digisF5HB.data,
-                             config.maxChannelsF5HB * sizeof(uint16_t) * compute_stride<Flavor5>(config.nsamplesF5HB)));
-        cudaCheck(cudaMalloc((void **)&digisF5HB.ids, sizeof(uint32_t) * config.maxChannelsF5HB));
-        cudaCheck(cudaMalloc((void **)&digisF5HB.npresamples, sizeof(uint8_t) * config.maxChannelsF5HB));
-
-        cudaCheck(cudaMalloc((void **)&digisF3HB.data,
-                             config.maxChannelsF3HB * sizeof(uint16_t) * compute_stride<Flavor3>(config.nsamplesF3HB)));
-        cudaCheck(cudaMalloc((void **)&digisF3HB.ids, config.maxChannelsF3HB * sizeof(uint32_t)));
-      }
-
-      void deallocate(ConfigurationParameters const &config) {
-        if (digisF01HE.data) {
-          cudaCheck(cudaFree(digisF01HE.data));
-          cudaCheck(cudaFree(digisF01HE.ids));
-        }
-
-        if (digisF5HB.data) {
-          cudaCheck(cudaFree(digisF5HB.data));
-          cudaCheck(cudaFree(digisF5HB.ids));
-          cudaCheck(cudaFree(digisF5HB.npresamples));
-        }
-
-        if (digisF3HB.data) {
-          cudaCheck(cudaFree(digisF3HB.data));
-          cudaCheck(cudaFree(digisF3HB.ids));
-        }
+      DigiCollection<Flavor01, common::DevStoragePolicy> digisF01HE;
+      DigiCollection<Flavor5, common::DevStoragePolicy> digisF5HB;
+      DigiCollection<Flavor3, common::DevStoragePolicy> digisF3HB;
+
+      void allocate(ConfigurationParameters const &config, cudaStream_t cudaStream) {
+        digisF01HE.data = cms::cuda::make_device_unique<uint16_t[]>(
+          config.maxChannelsF01HE*compute_stride<Flavor01>(config.nsamplesF01HE),
+          cudaStream
+        );
+        //cudaCheck(
+        //    cudaMalloc((void **)&digisF01HE.data,
+        //               config.maxChannelsF01HE * sizeof(uint16_t) * compute_stride<Flavor01>(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)));
       }
     };
 
     struct InputDataGPU {
-      unsigned char *data = nullptr;
-      uint32_t *offsets = nullptr;
-      int *feds = nullptr;
-
-      void allocate() {
-        cudaCheck(cudaMalloc((void **)&data, sizeof(unsigned char) * nbytes_per_fed_max * utca_nfeds_max));
-        cudaCheck(cudaMalloc((void **)&offsets, sizeof(uint32_t) * utca_nfeds_max));
-        cudaCheck(cudaMalloc((void **)&feds, sizeof(int) * utca_nfeds_max));
-      }
-
-      void deallocate() {
-        if (data) {
-          cudaCheck(cudaFree(data));
-          cudaCheck(cudaFree(offsets));
-          cudaCheck(cudaFree(feds));
-        }
-      }
+      cms::cuda::device::unique_ptr<unsigned char[]> data;
+      cms::cuda::device::unique_ptr<uint32_t[]> offsets;
+      cms::cuda::device::unique_ptr<int[]> feds;
     };
 
     struct ConditionsProducts {
diff --git a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
index 35cb128ba3d15..5011b7b74fac3 100644
--- a/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
+++ b/EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
@@ -568,29 +568,29 @@ namespace hcal {
                     uint32_t const nbytesTotal) {
       // transfer
       cudaCheck(cudaMemcpyAsync(
-          inputGPU.data, inputCPU.data.get(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
-      cudaCheck(cudaMemcpyAsync(inputGPU.offsets,
+          inputGPU.data.get(), inputCPU.data.get(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
+      cudaCheck(cudaMemcpyAsync(inputGPU.offsets.get(),
                                 inputCPU.offsets.get(),
                                 nfedsWithData * sizeof(uint32_t),
                                 cudaMemcpyHostToDevice,
                                 cudaStream));
       cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters.get(), 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
       cudaCheck(cudaMemcpyAsync(
-          inputGPU.feds, inputCPU.feds.get(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));
+          inputGPU.feds.get(), inputCPU.feds.get(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));
 
       // 12 is the max number of modules per crate
-      kernel_rawdecode_test<32><<<nfedsWithData, 12 * 32, 0, cudaStream>>>(inputGPU.data,
-                                                                           inputGPU.offsets,
-                                                                           inputGPU.feds,
+      kernel_rawdecode_test<32><<<nfedsWithData, 12 * 32, 0, cudaStream>>>(inputGPU.data.get(),
+                                                                           inputGPU.offsets.get(),
+                                                                           inputGPU.feds.get(),
                                                                            conditions.eMappingProduct.eid2did,
                                                                            conditions.eMappingProduct.eid2tid,
-                                                                           outputGPU.digisF01HE.data,
-                                                                           outputGPU.digisF01HE.ids,
-                                                                           outputGPU.digisF5HB.data,
-                                                                           outputGPU.digisF5HB.ids,
-                                                                           outputGPU.digisF5HB.npresamples,
-                                                                           outputGPU.digisF3HB.data,
-                                                                           outputGPU.digisF3HB.ids,
+                                                                           outputGPU.digisF01HE.data.get(),
+                                                                           outputGPU.digisF01HE.ids.get(),
+                                                                           outputGPU.digisF5HB.data.get(),
+                                                                           outputGPU.digisF5HB.ids.get(),
+                                                                           outputGPU.digisF5HB.npresamples.get(),
+                                                                           outputGPU.digisF3HB.data.get(),
+                                                                           outputGPU.digisF3HB.ids.get(),
                                                                            scratchGPU.pChannelsCounters.get(),
                                                                            config.nsamplesF01HE,
                                                                            config.nsamplesF5HB,
@@ -598,7 +598,7 @@ namespace hcal {
                                                                            nbytesTotal);
       cudaCheck(cudaGetLastError());
 
-      cudaCheck(cudaMemcpyAsync(outputCPU.nchannels.data(),
+      cudaCheck(cudaMemcpyAsync(outputCPU.nchannels.get(),
                                 scratchGPU.pChannelsCounters.get(),
                                 sizeof(uint32_t) * numOutputCollections,
                                 cudaMemcpyDeviceToHost,
diff --git a/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc b/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc
index 8aad10228021c..2a2fadf42ba12 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc
@@ -23,11 +23,11 @@ class HcalCPUDigisProducer : public edm::stream::EDProducer<edm::ExternalWork> {
   void produce(edm::Event&, edm::EventSetup const&) override;
 
 private:
-  using IProductTypef01 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::ViewStoragePolicy>>;
+  using IProductTypef01 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef01> digisF01HETokenIn_;
-  using IProductTypef5 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>;
+  using IProductTypef5 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef5> digisF5HBTokenIn_;
-  using IProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::ViewStoragePolicy>>;
+  using IProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef3> digisF3HBTokenIn_;
 
   using OProductTypef01 =
@@ -89,92 +89,27 @@ void HcalCPUDigisProducer::acquire(edm::Event const& event,
   digisf5HB_.resize(f5HBDigis.size);
   digisf3HB_.resize(f3HBDigis.size);
 
-  /*
-    idsf01he.resize(f01HEDigis.ndigis);
-    dataf01he.resize(f01HEDigis.ndigis * f01HEDigis.stride);
-    idsf5hb.resize(f5HBDigis.ndigis);
-    npresamplesf5hb.resize(f5HBDigis.ndigis);
-    dataf5hb.resize(f5HBDigis.ndigis * f5HBDigis.stride);
-    stridef01he = f01HEDigis.stride;
-    stridef5hb = f5HBDigis.stride;
-    */
-
   auto lambdaToTransfer = [&ctx](auto& dest, auto* src) {
     using vector_type = typename std::remove_reference<decltype(dest)>::type;
     using type = typename vector_type::value_type;
     cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream()));
   };
 
-  lambdaToTransfer(digisf01HE_.data, f01HEDigis.data);
-  lambdaToTransfer(digisf01HE_.ids, f01HEDigis.ids);
-
-  lambdaToTransfer(digisf5HB_.data, f5HBDigis.data);
-  lambdaToTransfer(digisf5HB_.ids, f5HBDigis.ids);
-  lambdaToTransfer(digisf5HB_.npresamples, f5HBDigis.npresamples);
-
-  lambdaToTransfer(digisf3HB_.data, f3HBDigis.data);
-  lambdaToTransfer(digisf3HB_.ids, f3HBDigis.ids);
-
-  /*
-    // enqeue transfers
-    cudaCheck( cudaMemcpyAsync(digisf01.data.data(),
-                               f01HEDigis.data,
-                               dataf01HE.data.size() * sizeof(uint16_t),
-                               cudaMemcpyDeviceToHost,
-                               ctx.stream().id()) );
-    cudaCheck( cudaMemcpyAsync(dataf5hb.data(),
-                               f5HBDigis.data,
-                               dataf5hb.size() * sizeof(uint16_t),
-                               cudaMemcpyDeviceToHost,
-                               ctx.stream().id()) );
-    cudaCheck( cudaMemcpyAsync(idsf01he.data(),
-                               f01HEDigis.ids,
-                               idsf01he.size() * sizeof(uint32_t),
-                               cudaMemcpyDeviceToHost,
-                               ctx.stream().id()) );
-    cudaCheck( cudaMemcpyAsync(idsf5hb.data(),
-                               f5HBDigis.ids,
-                               idsf5hb.size() * sizeof(uint32_t),
-                               cudaMemcpyDeviceToHost,
-                               ctx.stream().id()) );
-    cudaCheck( cudaMemcpyAsync(npresamplesf5hb.data(),
-                               f5HBDigis.npresamples,
-                               npresamplesf5hb.size() * sizeof(uint8_t),
-                               cudaMemcpyDeviceToHost,
-                               ctx.stream.id()) );
-                               */
+  lambdaToTransfer(digisf01HE_.data, f01HEDigis.data.get());
+  lambdaToTransfer(digisf01HE_.ids, f01HEDigis.ids.get());
+
+  lambdaToTransfer(digisf5HB_.data, f5HBDigis.data.get());
+  lambdaToTransfer(digisf5HB_.ids, f5HBDigis.ids.get());
+  lambdaToTransfer(digisf5HB_.npresamples, f5HBDigis.npresamples.get());
+
+  lambdaToTransfer(digisf3HB_.data, f3HBDigis.data.get());
+  lambdaToTransfer(digisf3HB_.ids, f3HBDigis.ids.get());
 }
 
 void HcalCPUDigisProducer::produce(edm::Event& event, edm::EventSetup const& setup) {
   event.emplace(digisF01HETokenOut_, std::move(digisf01HE_));
   event.emplace(digisF5HBTokenOut_, std::move(digisf5HB_));
   event.emplace(digisF3HBTokenOut_, std::move(digisf3HB_));
-
-  // output collections
-  /*
-    auto f01he = std::make_unique<edm::DataFrameContainer>(
-        stridef01he, HcalEndcap, idsf01he.size());
-    auto f5hb = std::make_unique<edm::DataFrameContainer>(
-        stridef5hb, HcalBarrel, idsf5hb.size());
-    
-    // cast constness away
-    // use pointers to buffers instead of move operator= semantics (or swap)
-    // cause we have different allocators in there...
-    auto *dataf01hetmp = const_cast<uint16_t*>(f01he->data().data());
-    auto *dataf5hbtmp = const_cast<uint16_t*>(f5hb->data().data());
-
-    auto *idsf01hetmp = const_cast<uint32_t*>(f01he->ids().data());
-    auto idsf5hbtmp = const_cast<uint32_t*>(f5hb->ids().data());
-
-    // copy data
-    std::memcpy(dataf01hetmp, dataf01he.data(), dataf01he.size() * sizeof(uint16_t));
-    std::memcpy(dataf5hbtmp, dataf5hb.data(), dataf5hb.size() * sizeof(uint16_t));
-    std::memcpy(idsf01hetmp, idsf01he.data(), idsf01he.size() * sizeof(uint32_t));
-    std::memcpy(idsf5hbtmp, idsf5hb.data(), idsf5hb.size() * sizeof(uint32_t));
-
-    event.put(digisF01HETokenOut_, std::move(f01he));
-    event.put(digisF5HBTokenOut_, std::move(f5hb));
-    */
 }
 
 DEFINE_FWK_MODULE(HcalCPUDigisProducer);
diff --git a/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc b/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
index c7bb27b60fa12..131f12d41f5c8 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
@@ -30,13 +30,13 @@ class HcalDigisProducerGPU : public edm::stream::EDProducer<edm::ExternalWork> {
   // type aliases
   using HostCollectionf01 =
       hcal::DigiCollection<hcal::Flavor01, hcal::common::VecStoragePolicy<hcal::CUDAHostAllocatorAlias>>;
-  using DeviceCollectionf01 = hcal::DigiCollection<hcal::Flavor01, hcal::common::ViewStoragePolicy>;
+  using DeviceCollectionf01 = hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>;
   using HostCollectionf5 =
       hcal::DigiCollection<hcal::Flavor5, hcal::common::VecStoragePolicy<hcal::CUDAHostAllocatorAlias>>;
-  using DeviceCollectionf5 = hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>;
+  using DeviceCollectionf5 = hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>;
   using HostCollectionf3 =
       hcal::DigiCollection<hcal::Flavor3, hcal::common::VecStoragePolicy<hcal::CUDAHostAllocatorAlias>>;
-  using DeviceCollectionf3 = hcal::DigiCollection<hcal::Flavor3, hcal::common::ViewStoragePolicy>;
+  using DeviceCollectionf3 = hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>;
 
   // output product tokens
   using ProductTypef01 = cms::cuda::Product<DeviceCollectionf01>;
@@ -48,16 +48,6 @@ class HcalDigisProducerGPU : public edm::stream::EDProducer<edm::ExternalWork> {
 
   cms::cuda::ContextState cudaState_;
 
-  /*
-    hcal::raw::ConfigurationParameters config_;
-    // FIXME move this to use raii
-    hcal::raw::InputDataCPU inputCPU_;
-    hcal::raw::InputDataGPU inputGPU_;
-    hcal::raw::OutputDataGPU outputGPU_;
-    hcal::raw::ScratchDataGPU scratchGPU_;
-    hcal::raw::OutputDataCPU outputCPU_;
-    */
-
   struct ConfigParameters {
     uint32_t maxChannelsF01HE, maxChannelsF5HB, maxChannelsF3HB, nsamplesF01HE, nsamplesF5HB, nsamplesF3HB;
   };
@@ -68,8 +58,7 @@ class HcalDigisProducerGPU : public edm::stream::EDProducer<edm::ExternalWork> {
   HostCollectionf5 hf5_;
   HostCollectionf3 hf3_;
 
-  // device products
-  // NOTE: this module owns memory of the product on the device
+  // device products: product owns memory (i.e. not the module)
   DeviceCollectionf01 df01_;
   DeviceCollectionf5 df5_;
   DeviceCollectionf3 df3_;
@@ -107,27 +96,6 @@ HcalDigisProducerGPU::HcalDigisProducerGPU(const edm::ParameterSet& ps)
   config_.nsamplesF5HB = ps.getParameter<uint32_t>("nsamplesF5HB");
   config_.nsamplesF3HB = ps.getParameter<uint32_t>("nsamplesF3HB");
 
-  // call CUDA API functions only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
-    // allocate on the device
-    cudaCheck(cudaMalloc(
-        (void**)&df01_.data,
-        config_.maxChannelsF01HE * sizeof(uint16_t) * hcal::compute_stride<hcal::Flavor01>(config_.nsamplesF01HE)));
-    cudaCheck(cudaMalloc((void**)&df01_.ids, config_.maxChannelsF01HE * sizeof(uint32_t)));
-
-    cudaCheck(cudaMalloc(
-        (void**)&df5_.data,
-        config_.maxChannelsF5HB * sizeof(uint16_t) * hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB)));
-    cudaCheck(cudaMalloc((void**)&df5_.ids, config_.maxChannelsF5HB * sizeof(uint32_t)));
-    cudaCheck(cudaMalloc((void**)&df5_.npresamples, sizeof(uint8_t) * config_.maxChannelsF5HB));
-
-    cudaCheck(cudaMalloc(
-        (void**)&df3_.data,
-        config_.maxChannelsF3HB * sizeof(uint16_t) * hcal::compute_stride<hcal::Flavor3>(config_.nsamplesF3HB)));
-    cudaCheck(cudaMalloc((void**)&df3_.ids, config_.maxChannelsF3HB * sizeof(uint32_t)));
-  }
-
   // preallocate on the host
   hf01_.stride = hcal::compute_stride<hcal::Flavor01>(config_.nsamplesF01HE);
   hf5_.stride = hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB);
@@ -137,22 +105,7 @@ HcalDigisProducerGPU::HcalDigisProducerGPU(const edm::ParameterSet& ps)
   hf3_.reserve(config_.maxChannelsF3HB);
 }
 
-HcalDigisProducerGPU::~HcalDigisProducerGPU() {
-  // call CUDA API functions only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
-    // deallocate on the device
-    cudaCheck(cudaFree(df01_.data));
-    cudaCheck(cudaFree(df01_.ids));
-
-    cudaCheck(cudaFree(df5_.data));
-    cudaCheck(cudaFree(df5_.ids));
-    cudaCheck(cudaFree(df5_.npresamples));
-
-    cudaCheck(cudaFree(df3_.data));
-    cudaCheck(cudaFree(df3_.ids));
-  }
-}
+HcalDigisProducerGPU::~HcalDigisProducerGPU() {}
 
 void HcalDigisProducerGPU::acquire(edm::Event const& event,
                                    edm::EventSetup const& setup,
@@ -170,6 +123,43 @@ void HcalDigisProducerGPU::acquire(edm::Event const& event,
   event.getByToken(hbheDigiToken_, hbheDigis);
   event.getByToken(qie11DigiToken_, qie11Digis);
 
+  // flavor 0/1 get devie blobs
+  df01_.data = cms::cuda::make_device_unique<uint16_t[]>(
+    config_.maxChannelsF01HE * hcal::compute_stride<hcal::Flavor01>(
+      config_.nsamplesF01HE),
+    ctx.stream()
+  );
+  df01_.ids = cms::cuda::make_device_unique<uint32_t[]>(
+    config_.maxChannelsF01HE,
+    ctx.stream()
+  );
+
+  // flavor3 get device blobs
+  df3_.data  = cms::cuda::make_device_unique<uint16_t[]>(
+    config_.maxChannelsF3HB * hcal::compute_stride<hcal::Flavor3>(
+      config_.nsamplesF3HB),
+    ctx.stream()
+  );
+  df3_.ids = cms::cuda::make_device_unique<uint32_t[]>(
+    config_.maxChannelsF3HB,
+    ctx.stream()
+  );
+
+  // flavor5 get device blobs
+  df5_.data = cms::cuda::make_device_unique<uint16_t[]>(
+    config_.maxChannelsF5HB * hcal::compute_stride<hcal::Flavor5>(
+      config_.nsamplesF5HB),
+    ctx.stream()
+  );
+  df5_.ids = cms::cuda::make_device_unique<uint32_t[]>(
+    config_.maxChannelsF5HB,
+    ctx.stream()
+  );
+  df5_.npresamples = cms::cuda::make_device_unique<uint8_t[]>(
+    config_.maxChannelsF5HB,
+    ctx.stream()
+  );
+
   for (auto const& hbhe : *hbheDigis) {
     auto const id = hbhe.id().rawId();
     auto const presamples = hbhe.presamples();
@@ -221,15 +211,15 @@ void HcalDigisProducerGPU::acquire(edm::Event const& event,
     cudaCheck(cudaMemcpyAsync(dest, src.data(), src.size() * sizeof(type), cudaMemcpyHostToDevice, ctx.stream()));
   };
 
-  lambdaToTransfer(df01_.data, hf01_.data);
-  lambdaToTransfer(df01_.ids, hf01_.ids);
+  lambdaToTransfer(df01_.data.get(), hf01_.data);
+  lambdaToTransfer(df01_.ids.get(), hf01_.ids);
 
-  lambdaToTransfer(df5_.data, hf5_.data);
-  lambdaToTransfer(df5_.ids, hf5_.ids);
-  lambdaToTransfer(df5_.npresamples, hf5_.npresamples);
+  lambdaToTransfer(df5_.data.get(), hf5_.data);
+  lambdaToTransfer(df5_.ids.get(), hf5_.ids);
+  lambdaToTransfer(df5_.npresamples.get(), hf5_.npresamples);
 
-  lambdaToTransfer(df3_.data, hf3_.data);
-  lambdaToTransfer(df3_.ids, hf3_.ids);
+  lambdaToTransfer(df3_.data.get(), hf3_.data);
+  lambdaToTransfer(df3_.ids.get(), hf3_.ids);
 }
 
 void HcalDigisProducerGPU::produce(edm::Event& event, edm::EventSetup const& setup) {
@@ -242,39 +232,9 @@ void HcalDigisProducerGPU::produce(edm::Event& event, edm::EventSetup const& set
   df3_.stride = hcal::compute_stride<hcal::Flavor3>(config_.nsamplesF3HB);
   df3_.size = hf3_.ids.size();
 
-  ctx.emplace(event, digisF01HEToken_, df01_);
-  ctx.emplace(event, digisF5HBToken_, df5_);
-  ctx.emplace(event, digisF3HBToken_, df3_);
-
-  /*
-
-#ifdef HCAL_RAWDECODE_CPUDEBUG
-    printf("f01he channels = %u f5hb channesl = %u\n",
-        outputCPU_.nchannels[hcal::raw::OutputF01HE], 
-        outputCPU_.nchannels[hcal::raw::OutputF5HB]);
-#endif
-
-    // FIXME: use sizes of views directly for cuda mem cpy?
-    auto const nchannelsF01HE = outputCPU_.nchannels[hcal::raw::OutputF01HE];
-    auto const nchannelsF5HB = outputCPU_.nchannels[hcal::raw::OutputF5HB];
-    outputGPU_.digisF01HE.size = nchannelsF01HE;
-    outputGPU_.digisF5HB.size = nchannelsF5HB;
-    outputGPU_.digisF01HE.stride = 
-        hcal::compute_stride<hcal::Flavor01>(config_.nsamplesF01HE);
-    outputGPU_.digisF5HB.stride = 
-        hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB);
-
-    hcal::DigiCollection<hcal::Flavor01> digisF01HE{outputGPU_.idsF01HE,
-        outputGPU_.digisF01HE, nchannelsF01HE, 
-        hcal::compute_stride<hcal::Flavor01>(config_.nsamplesF01HE)};
-    hcal::DigiCollection<hcal::Flavor5> digisF5HB{outputGPU_.idsF5HB,
-        outputGPU_.digisF5HB, outputGPU_.npresamplesF5HB, nchannelsF5HB, 
-        hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB)};
-
-    ctx.emplace(event, digisF01HEToken_, std::move(outputGPU_.digisF01HE));
-    ctx.emplace(event, digisF5HBToken_, std::move(outputGPU_.digisF5HB));
-
-    */
+  ctx.emplace(event, digisF01HEToken_, std::move(df01_));
+  ctx.emplace(event, digisF5HBToken_, std::move(df5_));
+  ctx.emplace(event, digisF3HBToken_, std::move(df3_));
 }
 
 DEFINE_FWK_MODULE(HcalDigisProducerGPU);
diff --git a/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc b/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
index ce5e85659ebdc..fd3c70d8a8160 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc
@@ -30,11 +30,11 @@ class HcalRawToDigiGPU : public edm::stream::EDProducer<edm::ExternalWork> {
 
 private:
   edm::EDGetTokenT<FEDRawDataCollection> rawDataToken_;
-  using ProductTypef01 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::ViewStoragePolicy>>;
+  using ProductTypef01 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>;
   edm::EDPutTokenT<ProductTypef01> digisF01HEToken_;
-  using ProductTypef5 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>;
+  using ProductTypef5 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>;
   edm::EDPutTokenT<ProductTypef5> digisF5HBToken_;
-  using ProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::ViewStoragePolicy>>;
+  using ProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>;
   edm::EDPutTokenT<ProductTypef3> digisF3HBToken_;
 
   cms::cuda::ContextState cudaState_;
@@ -42,8 +42,6 @@ class HcalRawToDigiGPU : public edm::stream::EDProducer<edm::ExternalWork> {
   std::vector<int> fedsToUnpack_;
 
   hcal::raw::ConfigurationParameters config_;
-  // FIXME move this to use raii
-  hcal::raw::InputDataGPU inputGPU_;
   hcal::raw::OutputDataGPU outputGPU_;
   hcal::raw::OutputDataCPU outputCPU_;
 };
@@ -83,24 +81,9 @@ HcalRawToDigiGPU::HcalRawToDigiGPU(const edm::ParameterSet& ps)
   config_.nsamplesF01HE = ps.getParameter<uint32_t>("nsamplesF01HE");
   config_.nsamplesF5HB = ps.getParameter<uint32_t>("nsamplesF5HB");
   config_.nsamplesF3HB = ps.getParameter<uint32_t>("nsamplesF3HB");
-
-  // reserve memory and call CUDA API functions only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
-    outputCPU_.allocate();
-
-    inputGPU_.allocate();
-    outputGPU_.allocate(config_);
-  }
 }
 
 HcalRawToDigiGPU::~HcalRawToDigiGPU() {
-  // call CUDA API functions only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
-    inputGPU_.deallocate();
-    outputGPU_.deallocate(config_);
-  }
 }
 
 void HcalRawToDigiGPU::acquire(edm::Event const& event,
@@ -141,6 +124,30 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,
       ctx.stream())
   };
 
+  // input data gpu
+  hcal::raw::InputDataGPU inputGPU = {
+    cms::cuda::make_device_unique<unsigned char[]>(
+      hcal::raw::utca_nfeds_max * hcal::raw::nbytes_per_fed_max,
+      ctx.stream()),
+    cms::cuda::make_device_unique<uint32_t[]>(
+      hcal::raw::utca_nfeds_max,
+      ctx.stream()),
+    cms::cuda::make_device_unique<int[]>(
+      hcal::raw::utca_nfeds_max,
+      ctx.stream())
+  };
+
+  // output cpu
+  outputCPU_ = {
+    cms::cuda::make_host_unique<uint32_t[]>(
+      hcal::raw::numOutputCollections,
+      ctx.stream()
+    )
+  };
+
+  // output gpu
+  outputGPU_.allocate(config_, ctx.stream());
+
   // iterate over feds
   // TODO: another idea
   //   - loop over all feds to unpack and enqueue cuda memcpy
@@ -173,7 +180,7 @@ void HcalRawToDigiGPU::acquire(edm::Event const& event,
   }
 
   hcal::raw::entryPoint(inputCPU,
-                        inputGPU_,
+                        inputGPU,
                         outputGPU_,
                         scratchGPU,
                         outputCPU_,
@@ -204,18 +211,12 @@ void HcalRawToDigiGPU::produce(edm::Event& event, edm::EventSetup const& setup)
   outputGPU_.digisF5HB.stride = hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB);
   outputGPU_.digisF3HB.stride = hcal::compute_stride<hcal::Flavor3>(config_.nsamplesF3HB);
 
-  /*
-    hcal::DigiCollection<hcal::Flavor01> digisF01HE{outputGPU_.idsF01HE,
-        outputGPU_.digisF01HE, nchannelsF01HE, 
-        hcal::compute_stride<hcal::Flavor01>(config_.nsamplesF01HE)};
-    hcal::DigiCollection<hcal::Flavor5> digisF5HB{outputGPU_.idsF5HB,
-        outputGPU_.digisF5HB, outputGPU_.npresamplesF5HB, nchannelsF5HB, 
-        hcal::compute_stride<hcal::Flavor5>(config_.nsamplesF5HB)};
-        */
-
   ctx.emplace(event, digisF01HEToken_, std::move(outputGPU_.digisF01HE));
   ctx.emplace(event, digisF5HBToken_, std::move(outputGPU_.digisF5HB));
   ctx.emplace(event, digisF3HBToken_, std::move(outputGPU_.digisF3HB));
+  
+  // reset ptrs that are carried as members
+  outputCPU_.nchannels.reset();
 }
 
 DEFINE_FWK_MODULE(HcalRawToDigiGPU);

From d05f147c6472b2bae37401d05f4652c09eca307f Mon Sep 17 00:00:00 2001
From: Viktor Khristenko <vdkhristenko1991@gmail.com>
Date: Fri, 3 Jul 2020 12:05:21 +0200
Subject: [PATCH 4/6] hcal/mahi/gpu switch to new input and use allocator for
 scratch

---
 .../bin/makeHcalRecHitGpuValidationPlots.cpp  |  2 +-
 .../HcalRecProducers/src/DeclsForKernels.h    | 16 ++++--
 .../src/HBHERecHitProducerGPU.cc              | 31 ++++++++---
 RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu | 54 +++++++++----------
 4 files changed, 64 insertions(+), 39 deletions(-)

diff --git a/RecoLocalCalo/HcalRecProducers/bin/makeHcalRecHitGpuValidationPlots.cpp b/RecoLocalCalo/HcalRecProducers/bin/makeHcalRecHitGpuValidationPlots.cpp
index 5ef7861f43232..e77727011e32b 100644
--- a/RecoLocalCalo/HcalRecProducers/bin/makeHcalRecHitGpuValidationPlots.cpp
+++ b/RecoLocalCalo/HcalRecProducers/bin/makeHcalRecHitGpuValidationPlots.cpp
@@ -61,7 +61,7 @@ int main(int argc, char* argv[]) {
   // prep input
   TFile rfin{inFileName.c_str()};
   TTree* rt = (TTree*)rfin.Get("Events");
-  rt->SetBranchAddress("HBHERecHitsSorted_hcalCPURecHitsProducer_recHitsLegacyLabelOut_RECO.", &wgpu);
+  rt->SetBranchAddress("HBHERecHitsSorted_hcalCPURecHitsProducer_recHitsLegacyHBHE_RECO.", &wgpu);
   //    rt->SetBranchAddress("hcalCUDAHostAllocatorAliashcalcommonVecStoragePolicyhcalRecHitCollection_hcalCPURecHitsProducer_recHitsM0LabelOut_RECO.", &wgpu);
   rt->SetBranchAddress("HBHERecHitsSorted_hbheprereco__RECO.", &wcpu);
 
diff --git a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
index 42f1992bcf119..aabfdb4b7813b 100644
--- a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
+++ b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
@@ -25,6 +25,9 @@
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalSiPMParametersGPU.h"
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalTimeCorrsGPU.h"
 
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
 namespace hcal {
   namespace mahi {
 
@@ -95,6 +98,11 @@ namespace hcal {
     };
 
     struct ScratchDataGPU {
+      cms::cuda::device::unique_ptr<float[]> amplitudes, noiseTerms,
+          pulseMatrices, pulseMatricesM, pulseMatricesP;
+      cms::cuda::device::unique_ptr<int8_t[]> soiSamples;
+
+      /*
       float *amplitudes = nullptr, *noiseTerms = nullptr;
       float *pulseMatrices = nullptr, *pulseMatricesM = nullptr, *pulseMatricesP = nullptr;
       int8_t* soiSamples = nullptr;
@@ -121,13 +129,13 @@ namespace hcal {
           cudaCheck(cudaFree(pulseMatricesP));
           cudaCheck(cudaFree(soiSamples));
         }
-      }
+      }*/
     };
 
     struct InputDataGPU {
-      DigiCollection<Flavor01, common::ViewStoragePolicy> const& f01HEDigis;
-      DigiCollection<Flavor5, common::ViewStoragePolicy> const& f5HBDigis;
-      DigiCollection<Flavor3, common::ViewStoragePolicy> const& f3HBDigis;
+      DigiCollection<Flavor01, common::DevStoragePolicy> const& f01HEDigis;
+      DigiCollection<Flavor5, common::DevStoragePolicy> const& f5HBDigis;
+      DigiCollection<Flavor3, common::DevStoragePolicy> const& f3HBDigis;
     };
 
   }  // namespace mahi
diff --git a/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc b/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc
index c78200badce57..9a760d88eb922 100644
--- a/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc
+++ b/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc
@@ -49,13 +49,13 @@ class HBHERecHitProducerGPU : public edm::stream::EDProducer<edm::ExternalWork>
   void acquire(edm::Event const&, edm::EventSetup const&, edm::WaitingTaskWithArenaHolder) override;
   void produce(edm::Event&, edm::EventSetup const&) override;
 
-  using IProductTypef01 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::ViewStoragePolicy>>;
+  using IProductTypef01 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef01> digisTokenF01HE_;
 
-  using IProductTypef5 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>;
+  using IProductTypef5 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef5> digisTokenF5HB_;
 
-  using IProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::ViewStoragePolicy>>;
+  using IProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef3> digisTokenF3HB_;
 
   using RecHitType = hcal::RecHitCollection<hcal::common::ViewStoragePolicy>;
@@ -64,7 +64,6 @@ class HBHERecHitProducerGPU : public edm::stream::EDProducer<edm::ExternalWork>
 
   hcal::mahi::ConfigParameters configParameters_;
   hcal::mahi::OutputDataGPU outputGPU_;
-  hcal::mahi::ScratchDataGPU scratchGPU_;
   cms::cuda::ContextState cudaState_;
 };
 
@@ -105,7 +104,6 @@ HBHERecHitProducerGPU::HBHERecHitProducerGPU(edm::ParameterSet const& ps)
   edm::Service<CUDAService> cs;
   if (cs and cs->enabled()) {
     outputGPU_.allocate(configParameters_);
-    scratchGPU_.allocate(configParameters_);
 
     // FIXME: use default device and default stream
     cudaCheck(
@@ -122,7 +120,6 @@ HBHERecHitProducerGPU::~HBHERecHitProducerGPU() {
   edm::Service<CUDAService> cs;
   if (cs and cs->enabled()) {
     outputGPU_.deallocate(configParameters_);
-    scratchGPU_.deallocate(configParameters_);
 
     cudaCheck(cudaFree(configParameters_.pulseOffsetsDevice));
   }
@@ -258,7 +255,27 @@ void HBHERecHitProducerGPU::acquire(edm::Event const& event,
                                             recConstantsHandle.product(),
                                             pedestalsHandle->offsetForHashes()};
 
-  hcal::mahi::entryPoint(inputGPU, outputGPU_, conditions, scratchGPU_, configParameters_, ctx.stream());
+  hcal::mahi::ScratchDataGPU scratchGPU = {
+    cms::cuda::make_device_unique<float[]>(
+      configParameters_.maxChannels * configParameters_.maxTimeSamples,
+      ctx.stream()
+    ),
+    cms::cuda::make_device_unique<float[]>(
+      configParameters_.maxChannels * configParameters_.maxTimeSamples, ctx.stream()),
+    cms::cuda::make_device_unique<float[]>(
+      configParameters_.maxChannels*configParameters_.maxTimeSamples*configParameters_.maxTimeSamples, 
+      ctx.stream()),
+    cms::cuda::make_device_unique<float[]>(
+      configParameters_.maxChannels*configParameters_.maxTimeSamples*configParameters_.maxTimeSamples, 
+      ctx.stream()),
+    cms::cuda::make_device_unique<float[]>(
+      configParameters_.maxChannels*configParameters_.maxTimeSamples*configParameters_.maxTimeSamples, 
+      ctx.stream()),
+    cms::cuda::make_device_unique<int8_t[]>(
+      configParameters_.maxChannels, ctx.stream()),
+  };
+
+  hcal::mahi::entryPoint(inputGPU, outputGPU_, conditions, scratchGPU, configParameters_, ctx.stream());
 
 #ifdef HCAL_MAHI_CPUDEBUG
   auto end = std::chrono::high_resolution_clock::now();
diff --git a/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu b/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
index 72df5d89815a2..8fcb5af45935e 100644
--- a/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
+++ b/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
@@ -1675,23 +1675,23 @@ namespace hcal {
       int nbytesShared =
           ((2 * f01nsamples + 2) * sizeof(float) + sizeof(uint64_t)) * configParameters.kprep1dChannelsPerBlock;
       kernel_prep1d_sameNumberOfSamples<<<blocks, threadsPerBlock, nbytesShared, cudaStream>>>(
-          scratch.amplitudes,
-          scratch.noiseTerms,
+          scratch.amplitudes.get(),
+          scratch.noiseTerms.get(),
           outputGPU.recHits.energy,
           outputGPU.recHits.chi2,
-          inputGPU.f01HEDigis.data,
-          inputGPU.f5HBDigis.data,
-          inputGPU.f3HBDigis.data,
-          inputGPU.f01HEDigis.ids,
-          inputGPU.f5HBDigis.ids,
-          inputGPU.f3HBDigis.ids,
+          inputGPU.f01HEDigis.data.get(),
+          inputGPU.f5HBDigis.data.get(),
+          inputGPU.f3HBDigis.data.get(),
+          inputGPU.f01HEDigis.ids.get(),
+          inputGPU.f5HBDigis.ids.get(),
+          inputGPU.f3HBDigis.ids.get(),
           inputGPU.f01HEDigis.stride,
           inputGPU.f5HBDigis.stride,
           inputGPU.f3HBDigis.stride,
           inputGPU.f01HEDigis.size,
           inputGPU.f5HBDigis.size,
-          inputGPU.f5HBDigis.npresamples,
-          scratch.soiSamples,
+          inputGPU.f5HBDigis.npresamples.get(),
+          scratch.soiSamples.get(),
           outputGPU.recHits.energyM0,
           outputGPU.recHits.timeM0,
           outputGPU.recHits.did,
@@ -1746,18 +1746,18 @@ namespace hcal {
 #endif
 
       kernel_prep_pulseMatrices_sameNumberOfSamples<<<blocks2, threadsPerBlock2, 0, cudaStream>>>(
-          scratch.pulseMatrices,
-          scratch.pulseMatricesM,
-          scratch.pulseMatricesP,
+          scratch.pulseMatrices.get(),
+          scratch.pulseMatricesM.get(),
+          scratch.pulseMatricesP.get(),
           configParameters.pulseOffsetsDevice,
-          scratch.amplitudes,
-          inputGPU.f01HEDigis.ids,
-          inputGPU.f5HBDigis.ids,
-          inputGPU.f3HBDigis.ids,
+          scratch.amplitudes.get(),
+          inputGPU.f01HEDigis.ids.get(),
+          inputGPU.f5HBDigis.ids.get(),
+          inputGPU.f3HBDigis.ids.get(),
           inputGPU.f01HEDigis.size,
           inputGPU.f5HBDigis.size,
           totalChannels,
-          scratch.soiSamples,
+          scratch.soiSamples.get(),
           conditions.recoParams.ids,
           conditions.recoParams.acc25nsVec,
           conditions.recoParams.diff25nsItvlVec,
@@ -1794,19 +1794,19 @@ namespace hcal {
         kernel_minimize<8, 8><<<blocks, threadsPerBlock, nbytesShared, cudaStream>>>(
             outputGPU.recHits.energy,
             outputGPU.recHits.chi2,
-            scratch.amplitudes,
-            scratch.pulseMatrices,
-            scratch.pulseMatricesM,
-            scratch.pulseMatricesP,
+            scratch.amplitudes.get(),
+            scratch.pulseMatrices.get(),
+            scratch.pulseMatricesM.get(),
+            scratch.pulseMatricesP.get(),
             configParameters.pulseOffsetsDevice,
-            scratch.noiseTerms,
-            scratch.soiSamples,
+            scratch.noiseTerms.get(),
+            scratch.soiSamples.get(),
             conditions.pedestalWidths.values,
             conditions.effectivePedestalWidths.values,
             configParameters.useEffectivePedestals,
-            inputGPU.f01HEDigis.ids,
-            inputGPU.f5HBDigis.ids,
-            inputGPU.f3HBDigis.ids,
+            inputGPU.f01HEDigis.ids.get(),
+            inputGPU.f5HBDigis.ids.get(),
+            inputGPU.f3HBDigis.ids.get(),
             conditions.gains.values,
             conditions.respCorrs.values,
             inputGPU.f01HEDigis.size,

From 8ee109b310238c3717840b56c0d7ca1bdf1af5d4 Mon Sep 17 00:00:00 2001
From: Viktor Khristenko <vdkhristenko1991@gmail.com>
Date: Fri, 3 Jul 2020 16:41:09 +0200
Subject: [PATCH 5/6] hcal/mahi/gpu move pulse offsets to use esproducer and
 all dev memory is relinquishable

---
 .../plugins/HcalCPUDigisProducer.cc           |  2 +
 .../plugins/HcalDigisProducerGPU.cc           |  2 +
 .../interface/HcalMahiPulseOffsetsGPU.h       | 39 +++++++++++
 .../src/HcalMahiPulseOffsetsGPU.cc            | 36 ++++++++++
 .../HcalRecProducers/src/DeclsForKernels.h    | 35 +++++-----
 .../src/HBHERecHitProducerGPU.cc              | 43 +++++-------
 .../src/HcalCPURecHitsProducer.cc             | 14 ++--
 .../src/HcalMahiPulseOffsetsGPUESProducer.cc  | 67 +++++++++++++++++++
 .../src/HcalMahiPulseOffsetsGPURecord.cc      |  4 ++
 .../src/HcalMahiPulseOffsetsGPURecord.h       |  7 ++
 RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu | 26 +++----
 11 files changed, 212 insertions(+), 63 deletions(-)
 create mode 100644 RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h
 create mode 100644 RecoLocalCalo/HcalRecAlgos/src/HcalMahiPulseOffsetsGPU.cc
 create mode 100644 RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPUESProducer.cc
 create mode 100644 RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.cc
 create mode 100644 RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.h

diff --git a/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc b/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc
index 2a2fadf42ba12..496714376367b 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalCPUDigisProducer.cc
@@ -92,6 +92,8 @@ void HcalCPUDigisProducer::acquire(edm::Event const& event,
   auto lambdaToTransfer = [&ctx](auto& dest, auto* src) {
     using vector_type = typename std::remove_reference<decltype(dest)>::type;
     using type = typename vector_type::value_type;
+    using src_data_type = typename std::remove_pointer<decltype(src)>::type;
+    static_assert(std::is_same<src_data_type, type>::value && "Dest and Src data types do not match");
     cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream()));
   };
 
diff --git a/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc b/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
index 131f12d41f5c8..944d1d249cd28 100644
--- a/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
+++ b/EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc
@@ -208,6 +208,8 @@ void HcalDigisProducerGPU::acquire(edm::Event const& event,
   auto lambdaToTransfer = [&ctx](auto* dest, auto const& src) {
     using vector_type = typename std::remove_reference<decltype(src)>::type;
     using type = typename vector_type::value_type;
+    using dest_data_type = typename std::remove_pointer<decltype(dest)>::type;
+    static_assert(std::is_same<dest_data_type, type>::value && "Dest and Src data typesdo not match");
     cudaCheck(cudaMemcpyAsync(dest, src.data(), src.size() * sizeof(type), cudaMemcpyHostToDevice, ctx.stream()));
   };
 
diff --git a/RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h b/RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h
new file mode 100644
index 0000000000000..2481a80711d33
--- /dev/null
+++ b/RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h
@@ -0,0 +1,39 @@
+#ifndef RecoLocalCalo_HcalRecAlgos_interface_HcalMahiPulseOffsetsGPU_h
+#define RecoLocalCalo_HcalRecAlgos_interface_HcalMahiPulseOffsetsGPU_h
+
+#include "FWCore/ParameterSet/interface/ParameterSet.h"
+
+#ifndef __CUDACC__
+#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
+#include "HeterogeneousCore/CUDACore/interface/ESProduct.h"
+#endif
+
+class HcalMahiPulseOffsetsGPU {
+public:
+  struct Product {
+    ~Product();
+    int* values;
+  };
+
+#ifndef __CUDACC__
+  // rearrange reco params
+  HcalMahiPulseOffsetsGPU(edm::ParameterSet const&);
+
+  // will trigger deallocation of Product thru ~Product
+  ~HcalMahiPulseOffsetsGPU() = default;
+
+  std::vector<int, cms::cuda::HostAllocator<int>> const& getValues() const {
+      return values_;
+  }
+
+  // get device pointers
+  Product const& getProduct(cudaStream_t) const;
+
+private:
+  std::vector<int, cms::cuda::HostAllocator<int>> values_;
+
+  cms::cuda::ESProduct<Product> product_;
+#endif
+};
+
+#endif
diff --git a/RecoLocalCalo/HcalRecAlgos/src/HcalMahiPulseOffsetsGPU.cc b/RecoLocalCalo/HcalRecAlgos/src/HcalMahiPulseOffsetsGPU.cc
new file mode 100644
index 0000000000000..3f5cdbe5f15ca
--- /dev/null
+++ b/RecoLocalCalo/HcalRecAlgos/src/HcalMahiPulseOffsetsGPU.cc
@@ -0,0 +1,36 @@
+#include "RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h"
+
+#include "FWCore/Utilities/interface/typelookup.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+
+// FIXME: add proper getters to conditions
+HcalMahiPulseOffsetsGPU::HcalMahiPulseOffsetsGPU(edm::ParameterSet const& ps) 
+{
+    auto const& values = ps.getParameter<std::vector<int>>("pulseOffsets");
+    values_.resize(values.size());
+    std::copy(values.begin(), values.end(), values_.begin());
+}
+
+HcalMahiPulseOffsetsGPU::Product::~Product() {
+  // deallocation
+  cudaCheck(cudaFree(values));
+}
+
+HcalMahiPulseOffsetsGPU::Product const& HcalMahiPulseOffsetsGPU::getProduct(cudaStream_t cudaStream) const {
+  auto const& product =
+      product_.dataForCurrentDeviceAsync(cudaStream, [this](HcalMahiPulseOffsetsGPU::Product& product, cudaStream_t cudaStream) {
+        // malloc
+        cudaCheck(cudaMalloc((void**)&product.values, this->values_.size() * sizeof(int)));
+
+        // transfer
+        cudaCheck(cudaMemcpyAsync(product.values,
+                                  this->values_.data(),
+                                  this->values_.size() * sizeof(int),
+                                  cudaMemcpyHostToDevice,
+                                  cudaStream));
+      });
+
+  return product;
+}
+
+TYPELOOKUP_DATA_REG(HcalMahiPulseOffsetsGPU);
diff --git a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
index aabfdb4b7813b..cce00a4e34c89 100644
--- a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
+++ b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
@@ -24,6 +24,7 @@
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalSiPMCharacteristicsGPU.h"
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalSiPMParametersGPU.h"
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalTimeCorrsGPU.h"
+#include "RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h"
 
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
@@ -49,6 +50,8 @@ namespace hcal {
       HcalTopology const* topology;
       HcalDDDRecConstants const* recConstants;
       uint32_t offsetForHashes;
+      HcalMahiPulseOffsetsGPU::Product const& pulseOffsets;
+      std::vector<int, cms::cuda::HostAllocator<int>> const& pulseOffsetsHost;
     };
 
     struct ConfigParameters {
@@ -65,7 +68,8 @@ namespace hcal {
       float ts4Thresh;
 
       std::vector<int> pulseOffsets;
-      int* pulseOffsetsDevice = nullptr;
+      // FIXME remove pulseOffsets - they come from esproduce now
+      //int* pulseOffsetsDevice = nullptr;
 
       std::array<uint32_t, 3> kernelMinimizeThreads;
 
@@ -78,22 +82,19 @@ namespace hcal {
     };
 
     struct OutputDataGPU {
-      RecHitCollection<common::ViewStoragePolicy> recHits;
-
-      void allocate(ConfigParameters const& config) {
-        cudaCheck(cudaMalloc((void**)&recHits.energy, config.maxChannels * sizeof(float)));
-        cudaCheck(cudaMalloc((void**)&recHits.chi2, config.maxChannels * sizeof(float)));
-        cudaCheck(cudaMalloc((void**)&recHits.energyM0, config.maxChannels * sizeof(float)));
-        cudaCheck(cudaMalloc((void**)&recHits.timeM0, config.maxChannels * sizeof(float)));
-        cudaCheck(cudaMalloc((void**)&recHits.did, config.maxChannels * sizeof(uint32_t)));
-      }
-
-      void deallocate(ConfigParameters const& config) {
-        cudaCheck(cudaFree(recHits.energy));
-        cudaCheck(cudaFree(recHits.chi2));
-        cudaCheck(cudaFree(recHits.energyM0));
-        cudaCheck(cudaFree(recHits.timeM0));
-        cudaCheck(cudaFree(recHits.did));
+      RecHitCollection<common::DevStoragePolicy> recHits;
+
+      void allocate(ConfigParameters const& config, cudaStream_t cudaStream) {
+        recHits.energy = cms::cuda::make_device_unique<float[]>(
+          config.maxChannels, cudaStream);
+        recHits.chi2 = cms::cuda::make_device_unique<float[]>(
+          config.maxChannels, cudaStream);
+        recHits.energyM0 = cms::cuda::make_device_unique<float[]>(
+          config.maxChannels, cudaStream);
+        recHits.timeM0 = cms::cuda::make_device_unique<float[]>(
+          config.maxChannels, cudaStream);
+        recHits.did = cms::cuda::make_device_unique<uint32_t[]>(
+          config.maxChannels, cudaStream);
       }
     };
 
diff --git a/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc b/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc
index 9a760d88eb922..32f768e78d164 100644
--- a/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc
+++ b/RecoLocalCalo/HcalRecProducers/src/HBHERecHitProducerGPU.cc
@@ -37,6 +37,9 @@
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalSiPMParametersGPU.h"
 #include "RecoLocalCalo/HcalRecAlgos/interface/HcalTimeCorrsGPU.h"
 
+#include "RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h"
+#include "HcalMahiPulseOffsetsGPURecord.h"
+
 #include "MahiGPU.h"
 
 class HBHERecHitProducerGPU : public edm::stream::EDProducer<edm::ExternalWork> {
@@ -58,7 +61,7 @@ class HBHERecHitProducerGPU : public edm::stream::EDProducer<edm::ExternalWork>
   using IProductTypef3 = cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductTypef3> digisTokenF3HB_;
 
-  using RecHitType = hcal::RecHitCollection<hcal::common::ViewStoragePolicy>;
+  using RecHitType = hcal::RecHitCollection<hcal::common::DevStoragePolicy>;
   using OProductType = cms::cuda::Product<RecHitType>;
   edm::EDPutTokenT<OProductType> rechitsM0Token_;
 
@@ -74,7 +77,6 @@ HBHERecHitProducerGPU::HBHERecHitProducerGPU(edm::ParameterSet const& ps)
       rechitsM0Token_{produces<OProductType>(ps.getParameter<std::string>("recHitsLabelM0HBHE"))} {
   configParameters_.maxChannels = ps.getParameter<uint32_t>("maxChannels");
   configParameters_.maxTimeSamples = ps.getParameter<uint32_t>("maxTimeSamples");
-  configParameters_.pulseOffsets = ps.getParameter<std::vector<int>>("pulseOffsets");
   configParameters_.kprep1dChannelsPerBlock = ps.getParameter<uint32_t>("kprep1dChannelsPerBlock");
   configParameters_.sipmQTSShift = ps.getParameter<int>("sipmQTSShift");
   configParameters_.sipmQNTStoSum = ps.getParameter<int>("sipmQNTStoSum");
@@ -99,37 +101,14 @@ HBHERecHitProducerGPU::HBHERecHitProducerGPU(edm::ParameterSet const& ps)
   configParameters_.kernelMinimizeThreads[0] = threadsMinimize[0];
   configParameters_.kernelMinimizeThreads[1] = threadsMinimize[1];
   configParameters_.kernelMinimizeThreads[2] = threadsMinimize[2];
-
-  // call CUDA API functions only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
-    outputGPU_.allocate(configParameters_);
-
-    // FIXME: use default device and default stream
-    cudaCheck(
-        cudaMalloc((void**)&configParameters_.pulseOffsetsDevice, sizeof(int) * configParameters_.pulseOffsets.size()));
-    cudaCheck(cudaMemcpy(configParameters_.pulseOffsetsDevice,
-                         configParameters_.pulseOffsets.data(),
-                         configParameters_.pulseOffsets.size() * sizeof(int),
-                         cudaMemcpyHostToDevice));
-  }
 }
 
-HBHERecHitProducerGPU::~HBHERecHitProducerGPU() {
-  // call CUDA API functions only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
-    outputGPU_.deallocate(configParameters_);
-
-    cudaCheck(cudaFree(configParameters_.pulseOffsetsDevice));
-  }
-}
+HBHERecHitProducerGPU::~HBHERecHitProducerGPU() {}
 
 void HBHERecHitProducerGPU::fillDescriptions(edm::ConfigurationDescriptions& cdesc) {
   edm::ParameterSetDescription desc;
   desc.add<uint32_t>("maxChannels", 10000u);
   desc.add<uint32_t>("maxTimeSamples", 10);
-  desc.add<std::vector<int>>("pulseOffsets", {-3, -2, -1, 0, 1, 2, 3, 4});
   desc.add<uint32_t>("kprep1dChannelsPerBlock", 32);
   desc.add<edm::InputTag>("digisLabelF01HE", edm::InputTag{"hcalRawToDigiGPU", "f01HEDigisGPU"});
   desc.add<edm::InputTag>("digisLabelF5HB", edm::InputTag{"hcalRawToDigiGPU", "f5HBDigisGPU"});
@@ -236,6 +215,10 @@ void HBHERecHitProducerGPU::acquire(edm::Event const& event,
   setup.get<HcalSiPMCharacteristicsRcd>().get(sipmCharacteristicsHandle);
   auto const& sipmCharacteristicsProduct = sipmCharacteristicsHandle->getProduct(ctx.stream());
 
+  edm::ESHandle<HcalMahiPulseOffsetsGPU> pulseOffsetsHandle;
+  setup.get<HcalMahiPulseOffsetsGPURecord>().get(pulseOffsetsHandle);
+  auto const& pulseOffsetsProduct = pulseOffsetsHandle->getProduct(ctx.stream());
+
   // bundle up conditions
   hcal::mahi::ConditionsProducts conditions{gainWidthsProduct,
                                             gainsProduct,
@@ -253,8 +236,11 @@ void HBHERecHitProducerGPU::acquire(edm::Event const& event,
                                             effectivePedestalsProduct,
                                             topologyHandle.product(),
                                             recConstantsHandle.product(),
-                                            pedestalsHandle->offsetForHashes()};
+                                            pedestalsHandle->offsetForHashes(),
+                                            pulseOffsetsProduct,
+                                            pulseOffsetsHandle->getValues()};
 
+  // scratch mem on device
   hcal::mahi::ScratchDataGPU scratchGPU = {
     cms::cuda::make_device_unique<float[]>(
       configParameters_.maxChannels * configParameters_.maxTimeSamples,
@@ -275,6 +261,9 @@ void HBHERecHitProducerGPU::acquire(edm::Event const& event,
       configParameters_.maxChannels, ctx.stream()),
   };
 
+  // output dev mem
+  outputGPU_.allocate(configParameters_, ctx.stream());
+
   hcal::mahi::entryPoint(inputGPU, outputGPU_, conditions, scratchGPU, configParameters_, ctx.stream());
 
 #ifdef HCAL_MAHI_CPUDEBUG
diff --git a/RecoLocalCalo/HcalRecProducers/src/HcalCPURecHitsProducer.cc b/RecoLocalCalo/HcalRecProducers/src/HcalCPURecHitsProducer.cc
index db934710f6108..e46f9712651a9 100644
--- a/RecoLocalCalo/HcalRecProducers/src/HcalCPURecHitsProducer.cc
+++ b/RecoLocalCalo/HcalRecProducers/src/HcalCPURecHitsProducer.cc
@@ -23,7 +23,7 @@ class HcalCPURecHitsProducer : public edm::stream::EDProducer<edm::ExternalWork>
   void produce(edm::Event&, edm::EventSetup const&) override;
 
 private:
-  using IProductType = cms::cuda::Product<hcal::RecHitCollection<hcal::common::ViewStoragePolicy>>;
+  using IProductType = cms::cuda::Product<hcal::RecHitCollection<hcal::common::DevStoragePolicy>>;
   edm::EDGetTokenT<IProductType> recHitsM0TokenIn_;
   using OProductType = hcal::RecHitCollection<hcal::common::VecStoragePolicy<hcal::CUDAHostAllocatorAlias>>;
   edm::EDPutTokenT<OProductType> recHitsM0TokenOut_;
@@ -67,15 +67,17 @@ void HcalCPURecHitsProducer::acquire(edm::Event const& event,
 
   auto lambdaToTransfer = [&ctx](auto& dest, auto* src) {
     using vector_type = typename std::remove_reference<decltype(dest)>::type;
+    using src_data_type = typename std::remove_pointer<decltype(src)>::type;
     using type = typename vector_type::value_type;
+    static_assert(std::is_same<src_data_type, type>::value && "Dest and Src data types do not match");
     cudaCheck(cudaMemcpyAsync(dest.data(), src, dest.size() * sizeof(type), cudaMemcpyDeviceToHost, ctx.stream()));
   };
 
-  lambdaToTransfer(tmpRecHits_.energy, recHits.energy);
-  lambdaToTransfer(tmpRecHits_.chi2, recHits.chi2);
-  lambdaToTransfer(tmpRecHits_.energyM0, recHits.energyM0);
-  lambdaToTransfer(tmpRecHits_.timeM0, recHits.timeM0);
-  lambdaToTransfer(tmpRecHits_.did, recHits.did);
+  lambdaToTransfer(tmpRecHits_.energy, recHits.energy.get());
+  lambdaToTransfer(tmpRecHits_.chi2, recHits.chi2.get());
+  lambdaToTransfer(tmpRecHits_.energyM0, recHits.energyM0.get());
+  lambdaToTransfer(tmpRecHits_.timeM0, recHits.timeM0.get());
+  lambdaToTransfer(tmpRecHits_.did, recHits.did.get());
 }
 
 void HcalCPURecHitsProducer::produce(edm::Event& event, edm::EventSetup const& setup) {
diff --git a/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPUESProducer.cc b/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPUESProducer.cc
new file mode 100644
index 0000000000000..9500f62ea869f
--- /dev/null
+++ b/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPUESProducer.cc
@@ -0,0 +1,67 @@
+#include <array>
+#include <iostream>
+#include <tuple>
+#include <utility>
+
+#include "FWCore/Framework/interface/ESProducer.h"
+#include "FWCore/Framework/interface/ESProductHost.h"
+#include "FWCore/Framework/interface/ESTransientHandle.h"
+#include "FWCore/Framework/interface/MakerMacros.h"
+#include "FWCore/Framework/interface/ModuleFactory.h"
+#include "FWCore/Framework/interface/EventSetupRecordIntervalFinder.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+#include "FWCore/ParameterSet/interface/ParameterSet.h"
+#include "FWCore/Utilities/interface/ReusableObjectHolder.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+#include "RecoLocalCalo/HcalRecAlgos/interface/HcalMahiPulseOffsetsGPU.h"
+#include "HcalMahiPulseOffsetsGPURecord.h"
+
+#include "FWCore/Framework/interface/SourceFactory.h"
+
+class HcalMahiPulseOffsetsGPUESProducer 
+        : public edm::ESProducer, public edm::EventSetupRecordIntervalFinder {
+public:
+    HcalMahiPulseOffsetsGPUESProducer(edm::ParameterSet const&);
+    ~HcalMahiPulseOffsetsGPUESProducer() override = default;
+
+    static void fillDescriptions(edm::ConfigurationDescriptions&);
+    std::unique_ptr<HcalMahiPulseOffsetsGPU> produce(HcalMahiPulseOffsetsGPURecord const&);
+
+protected:
+    void setIntervalFor(const edm::eventsetup::EventSetupRecordKey&,
+                        const edm::IOVSyncValue&,
+                        edm::ValidityInterval&) override;
+
+private:
+    edm::ParameterSet const& pset_;
+};
+
+HcalMahiPulseOffsetsGPUESProducer::HcalMahiPulseOffsetsGPUESProducer(
+        edm::ParameterSet const& pset) : pset_{pset}
+{
+    setWhatProduced(this);
+    findingRecord<HcalMahiPulseOffsetsGPURecord>();
+}
+
+void HcalMahiPulseOffsetsGPUESProducer::setIntervalFor(
+        const edm::eventsetup::EventSetupRecordKey& iKey,
+        const edm::IOVSyncValue& iTime,
+        edm::ValidityInterval& oInterval) {
+    oInterval = edm::ValidityInterval(
+        edm::IOVSyncValue::beginOfTime(), edm::IOVSyncValue::endOfTime());
+}
+
+void HcalMahiPulseOffsetsGPUESProducer::fillDescriptions(
+        edm::ConfigurationDescriptions& desc) {
+    edm::ParameterSetDescription d;
+    d.add<std::vector<int>>("pulseOffsets", {-3, -2, -1, 0, 1, 2, 3, 4});
+    desc.addWithDefaultLabel(d);
+}
+
+std::unique_ptr<HcalMahiPulseOffsetsGPU> HcalMahiPulseOffsetsGPUESProducer::produce(
+        HcalMahiPulseOffsetsGPURecord const&) {
+    return std::make_unique<HcalMahiPulseOffsetsGPU>(pset_);
+}
+
+DEFINE_FWK_EVENTSETUP_SOURCE(HcalMahiPulseOffsetsGPUESProducer);
diff --git a/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.cc b/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.cc
new file mode 100644
index 0000000000000..e59e802411c9b
--- /dev/null
+++ b/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.cc
@@ -0,0 +1,4 @@
+#include "HcalMahiPulseOffsetsGPURecord.h"
+
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+EVENTSETUP_RECORD_REG(HcalMahiPulseOffsetsGPURecord);
diff --git a/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.h b/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.h
new file mode 100644
index 0000000000000..ea6683606ca72
--- /dev/null
+++ b/RecoLocalCalo/HcalRecProducers/src/HcalMahiPulseOffsetsGPURecord.h
@@ -0,0 +1,7 @@
+#ifndef HcalMahiPulseOffsetsGPURecord_h
+#define HcalMahiPulseOffsetsGPURecord_h
+
+#include "FWCore/Framework/interface/EventSetupRecordImplementation.h"
+class HcalMahiPulseOffsetsGPURecord : public edm::eventsetup::EventSetupRecordImplementation<HcalMahiPulseOffsetsGPURecord> {};
+
+#endif
diff --git a/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu b/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
index 8fcb5af45935e..05ddb83ae540a 100644
--- a/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
+++ b/RecoLocalCalo/HcalRecProducers/src/MahiGPU.cu
@@ -1677,8 +1677,8 @@ namespace hcal {
       kernel_prep1d_sameNumberOfSamples<<<blocks, threadsPerBlock, nbytesShared, cudaStream>>>(
           scratch.amplitudes.get(),
           scratch.noiseTerms.get(),
-          outputGPU.recHits.energy,
-          outputGPU.recHits.chi2,
+          outputGPU.recHits.energy.get(),
+          outputGPU.recHits.chi2.get(),
           inputGPU.f01HEDigis.data.get(),
           inputGPU.f5HBDigis.data.get(),
           inputGPU.f3HBDigis.data.get(),
@@ -1692,9 +1692,9 @@ namespace hcal {
           inputGPU.f5HBDigis.size,
           inputGPU.f5HBDigis.npresamples.get(),
           scratch.soiSamples.get(),
-          outputGPU.recHits.energyM0,
-          outputGPU.recHits.timeM0,
-          outputGPU.recHits.did,
+          outputGPU.recHits.energyM0.get(),
+          outputGPU.recHits.timeM0.get(),
+          outputGPU.recHits.did.get(),
           totalChannels,
           conditions.recoParams.param1,
           conditions.recoParams.param2,
@@ -1734,8 +1734,8 @@ namespace hcal {
 
       // 1024 is the max threads per block for gtx1080
       // FIXME: take this from cuda service or something like that
-      uint32_t const channelsPerBlock = 1024 / (f01nsamples * configParameters.pulseOffsets.size());
-      dim3 threadsPerBlock2{f01nsamples, static_cast<uint32_t>(configParameters.pulseOffsets.size()), channelsPerBlock};
+      uint32_t const channelsPerBlock = 1024 / (f01nsamples * conditions.pulseOffsetsHost.size());
+      dim3 threadsPerBlock2{f01nsamples, static_cast<uint32_t>(conditions.pulseOffsetsHost.size()), channelsPerBlock};
       int blocks2 =
           threadsPerBlock2.z > totalChannels ? 1 : (totalChannels + threadsPerBlock2.z - 1) / threadsPerBlock2.z;
 
@@ -1749,7 +1749,7 @@ namespace hcal {
           scratch.pulseMatrices.get(),
           scratch.pulseMatricesM.get(),
           scratch.pulseMatricesP.get(),
-          configParameters.pulseOffsetsDevice,
+          conditions.pulseOffsets.values,
           scratch.amplitudes.get(),
           inputGPU.f01HEDigis.ids.get(),
           inputGPU.f5HBDigis.ids.get(),
@@ -1786,19 +1786,19 @@ namespace hcal {
           configParameters.tmaxTimeSlew);
       cudaCheck(cudaGetLastError());
 
-      if (f01nsamples == 8 && configParameters.pulseOffsets.size() == 8u) {
+      if (f01nsamples == 8 && conditions.pulseOffsetsHost.size() == 8u) {
         // FIXME: provide constants from configuration
         uint32_t threadsPerBlock = configParameters.kernelMinimizeThreads[0];
         uint32_t blocks = threadsPerBlock > totalChannels ? 1 : (totalChannels + threadsPerBlock - 1) / threadsPerBlock;
         auto const nbytesShared = 2 * threadsPerBlock * MapSymM<float, 8>::total * sizeof(float);
         kernel_minimize<8, 8><<<blocks, threadsPerBlock, nbytesShared, cudaStream>>>(
-            outputGPU.recHits.energy,
-            outputGPU.recHits.chi2,
+            outputGPU.recHits.energy.get(),
+            outputGPU.recHits.chi2.get(),
             scratch.amplitudes.get(),
             scratch.pulseMatrices.get(),
             scratch.pulseMatricesM.get(),
             scratch.pulseMatricesP.get(),
-            configParameters.pulseOffsetsDevice,
+            conditions.pulseOffsets.values,
             scratch.noiseTerms.get(),
             scratch.soiSamples.get(),
             conditions.pedestalWidths.values,
@@ -1827,7 +1827,7 @@ namespace hcal {
       } else {
         throw cms::Exception("Invalid MahiGPU configuration")
             << "Currently support only 8 pulses and 8 time samples and provided: " << f01nsamples << " samples and "
-            << configParameters.pulseOffsets.size() << " pulses" << std::endl;
+            << conditions.pulseOffsetsHost.size() << " pulses" << std::endl;
       }
     }
 

From 9e67a225e37530224eb77dde8077674ae242ad82 Mon Sep 17 00:00:00 2001
From: Viktor Khristenko <vdkhristenko1991@gmail.com>
Date: Sat, 4 Jul 2020 13:16:18 +0200
Subject: [PATCH 6/6] remove old

---
 RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
index cce00a4e34c89..841e2b04f12fd 100644
--- a/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
+++ b/RecoLocalCalo/HcalRecProducers/src/DeclsForKernels.h
@@ -67,10 +67,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: