Synchronise with CMSSW_11_2_0_pre6

CUDADataFormats/CaloCommon/BuildFile.xml

@@ -0,0 +1,6 @@
+<use name="CUDADataFormats/Common" />
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+  <lib   name="1"/>
+</export>

CUDADataFormats/CaloCommon/interface/Common.h

@@ -0,0 +1,68 @@
+#ifndef CUDADataFormats_CaloCommon_interface_Common_h
+#define CUDADataFormats_CaloCommon_interface_Common_h
+
+#include <vector>
+
+#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+
+namespace calo {
+  namespace common {
+
+    // FIXME: not able to get enums to work with genreflex
+    namespace tags {
+
+      struct Vec {};
+      struct Ptr {};
+      struct DevPtr {};
+
+    }  // namespace tags
+
+    template <typename tag>
+    struct AddSize {};
+
+    template <>
+    struct AddSize<tags::Ptr> {
+      uint32_t size;
+    };
+
+    template <>
+    struct AddSize<tags::DevPtr> {
+      uint32_t size;
+    };
+
+    struct ViewStoragePolicy {
+      using TagType = tags::Ptr;
+
+      template <typename T>
+      struct StorageSelector {
+        using type = T*;
+      };
+    };
+
+    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;
+
+      template <typename T>
+      struct StorageSelector {
+        using type = std::vector<T, Allocator<T>>;
+      };
+    };
+
+    template <typename T>
+    using CUDAHostAllocatorAlias = cms::cuda::HostAllocator<T>;
+
+  }  // namespace common
+}  // namespace calo
+
+#endif  // CUDADataFormats_CaloCommon_interface_Common_h

CUDADataFormats/CaloCommon/src/classes.h

@@ -0,0 +1,2 @@
+#include "DataFormats/Common/interface/Wrapper.h"
+#include "CUDADataFormats/CaloCommon/interface/Common.h"

CUDADataFormats/CaloCommon/src/classes_def.xml

@@ -0,0 +1,8 @@
+<lcgdict>
+    <enum name="calo::common::tags::Ptr"/>
+    <enum name="calo::common::tags::Vec"/>
+    <enum name="calo::common::tags::DevPtr"/>
+    <class name="calo::common::AddSize<calo::common::tags::Ptr>"/>
+    <class name="calo::common::AddSize<calo::common::tags::Vec>"/>
+    <class name="calo::common::AddSize<calo::common::tags::DevPtr>"/>
+</lcgdict>

CUDADataFormats/Common/interface/ArrayShadow.h

@@ -0,0 +1,12 @@
+#ifndef CUDADataFormatsCommonArrayShadow_H
+#define CUDADataFormatsCommonArrayShadow_H
+#include <array>
+
+template <typename A>
+struct ArrayShadow {
+  using T = typename A::value_type;
+  constexpr static auto size() { return std::tuple_size<A>::value; }
+  T data[std::tuple_size<A>::value];
+};
+
+#endif

CUDADataFormats/Common/interface/HeterogeneousSoA.h

@@ -0,0 +1,189 @@
+#ifndef CUDADataFormatsCommonHeterogeneousSoA_H
+#define CUDADataFormatsCommonHeterogeneousSoA_H
+
+#include <cassert>
+
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+// a heterogeneous unique pointer...
+template <typename T>
+class HeterogeneousSoA {
+public:
+  using Product = T;
+
+  HeterogeneousSoA() = default;  // make root happy
+  ~HeterogeneousSoA() = default;
+  HeterogeneousSoA(HeterogeneousSoA &&) = default;
+  HeterogeneousSoA &operator=(HeterogeneousSoA &&) = default;
+
+  explicit HeterogeneousSoA(cms::cuda::device::unique_ptr<T> &&p) : dm_ptr(std::move(p)) {}
+  explicit HeterogeneousSoA(cms::cuda::host::unique_ptr<T> &&p) : hm_ptr(std::move(p)) {}
+  explicit HeterogeneousSoA(std::unique_ptr<T> &&p) : std_ptr(std::move(p)) {}
+
+  auto const *get() const { return dm_ptr ? dm_ptr.get() : (hm_ptr ? hm_ptr.get() : std_ptr.get()); }
+
+  auto const &operator*() const { return *get(); }
+
+  auto const *operator->() const { return get(); }
+
+  auto *get() { return dm_ptr ? dm_ptr.get() : (hm_ptr ? hm_ptr.get() : std_ptr.get()); }
+
+  auto &operator*() { return *get(); }
+
+  auto *operator->() { return get(); }
+
+  // in reality valid only for GPU version...
+  cms::cuda::host::unique_ptr<T> toHostAsync(cudaStream_t stream) const {
+    assert(dm_ptr);
+    auto ret = cms::cuda::make_host_unique<T>(stream);
+    cudaCheck(cudaMemcpyAsync(ret.get(), dm_ptr.get(), sizeof(T), cudaMemcpyDefault, stream));
+    return ret;
+  }
+
+private:
+  // a union wan't do it, a variant will not be more efficienct
+  cms::cuda::device::unique_ptr<T> dm_ptr;  //!
+  cms::cuda::host::unique_ptr<T> hm_ptr;    //!
+  std::unique_ptr<T> std_ptr;               //!
+};
+
+namespace cms {
+  namespace cudacompat {
+
+    struct GPUTraits {
+      template <typename T>
+      using unique_ptr = cms::cuda::device::unique_ptr<T>;
+
+      template <typename T>
+      static auto make_unique(cudaStream_t stream) {
+        return cms::cuda::make_device_unique<T>(stream);
+      }
+
+      template <typename T>
+      static auto make_unique(size_t size, cudaStream_t stream) {
+        return cms::cuda::make_device_unique<T>(size, stream);
+      }
+
+      template <typename T>
+      static auto make_host_unique(cudaStream_t stream) {
+        return cms::cuda::make_host_unique<T>(stream);
+      }
+
+      template <typename T>
+      static auto make_device_unique(cudaStream_t stream) {
+        return cms::cuda::make_device_unique<T>(stream);
+      }
+
+      template <typename T>
+      static auto make_device_unique(size_t size, cudaStream_t stream) {
+        return cms::cuda::make_device_unique<T>(size, stream);
+      }
+    };
+
+    struct HostTraits {
+      template <typename T>
+      using unique_ptr = cms::cuda::host::unique_ptr<T>;
+
+      template <typename T>
+      static auto make_unique(cudaStream_t stream) {
+        return cms::cuda::make_host_unique<T>(stream);
+      }
+
+      template <typename T>
+      static auto make_host_unique(cudaStream_t stream) {
+        return cms::cuda::make_host_unique<T>(stream);
+      }
+
+      template <typename T>
+      static auto make_device_unique(cudaStream_t stream) {
+        return cms::cuda::make_device_unique<T>(stream);
+      }
+
+      template <typename T>
+      static auto make_device_unique(size_t size, cudaStream_t stream) {
+        return cms::cuda::make_device_unique<T>(size, stream);
+      }
+    };
+
+    struct CPUTraits {
+      template <typename T>
+      using unique_ptr = std::unique_ptr<T>;
+
+      template <typename T>
+      static auto make_unique(cudaStream_t) {
+        return std::make_unique<T>();
+      }
+
+      template <typename T>
+      static auto make_unique(size_t size, cudaStream_t) {
+        return std::make_unique<T>(size);
+      }
+
+      template <typename T>
+      static auto make_host_unique(cudaStream_t) {
+        return std::make_unique<T>();
+      }
+
+      template <typename T>
+      static auto make_device_unique(cudaStream_t) {
+        return std::make_unique<T>();
+      }
+
+      template <typename T>
+      static auto make_device_unique(size_t size, cudaStream_t) {
+        return std::make_unique<T>(size);
+      }
+    };
+
+  }  // namespace cudacompat
+}  // namespace cms
+
+// a heterogeneous unique pointer (of a different sort) ...
+template <typename T, typename Traits>
+class HeterogeneousSoAImpl {
+public:
+  template <typename V>
+  using unique_ptr = typename Traits::template unique_ptr<V>;
+
+  HeterogeneousSoAImpl() = default;  // make root happy
+  ~HeterogeneousSoAImpl() = default;
+  HeterogeneousSoAImpl(HeterogeneousSoAImpl &&) = default;
+  HeterogeneousSoAImpl &operator=(HeterogeneousSoAImpl &&) = default;
+
+  explicit HeterogeneousSoAImpl(unique_ptr<T> &&p) : m_ptr(std::move(p)) {}
+  explicit HeterogeneousSoAImpl(cudaStream_t stream);
+
+  T const *get() const { return m_ptr.get(); }
+
+  T *get() { return m_ptr.get(); }
+
+  cms::cuda::host::unique_ptr<T> toHostAsync(cudaStream_t stream) const;
+
+private:
+  unique_ptr<T> m_ptr;  //!
+};
+
+template <typename T, typename Traits>
+HeterogeneousSoAImpl<T, Traits>::HeterogeneousSoAImpl(cudaStream_t stream) {
+  m_ptr = Traits::template make_unique<T>(stream);
+}
+
+// in reality valid only for GPU version...
+template <typename T, typename Traits>
+cms::cuda::host::unique_ptr<T> HeterogeneousSoAImpl<T, Traits>::toHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<T>(stream);
+  cudaCheck(cudaMemcpyAsync(ret.get(), get(), sizeof(T), cudaMemcpyDefault, stream));
+  return ret;
+}
+
+template <typename T>
+using HeterogeneousSoAGPU = HeterogeneousSoAImpl<T, cms::cudacompat::GPUTraits>;
+template <typename T>
+using HeterogeneousSoACPU = HeterogeneousSoAImpl<T, cms::cudacompat::CPUTraits>;
+template <typename T>
+using HeterogeneousSoAHost = HeterogeneousSoAImpl<T, cms::cudacompat::HostTraits>;
+
+#endif

CUDADataFormats/Common/interface/HostProduct.h

@@ -0,0 +1,29 @@
+#ifndef CUDADataFormatsCommonHostProduct_H
+#define CUDADataFormatsCommonHostProduct_H
+
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+// a heterogeneous unique pointer...
+template <typename T>
+class HostProduct {
+public:
+  HostProduct() = default;  // make root happy
+  ~HostProduct() = default;
+  HostProduct(HostProduct&&) = default;
+  HostProduct& operator=(HostProduct&&) = default;
+
+  explicit HostProduct(cms::cuda::host::unique_ptr<T>&& p) : hm_ptr(std::move(p)) {}
+  explicit HostProduct(std::unique_ptr<T>&& p) : std_ptr(std::move(p)) {}
+
+  auto const* get() const { return hm_ptr ? hm_ptr.get() : std_ptr.get(); }
+
+  auto const& operator*() const { return *get(); }
+
+  auto const* operator->() const { return get(); }
+
+private:
+  cms::cuda::host::unique_ptr<T> hm_ptr;  //!
+  std::unique_ptr<T> std_ptr;             //!
+};
+
+#endif

CUDADataFormats/EcalDigi/BuildFile.xml

@@ -0,0 +1,8 @@
+<use name="CUDADataFormats/Common"/>
+<use name="CUDADataFormats/CaloCommon"/>
+<use name="DataFormats/Common"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+  <lib   name="1"/>
+</export>

CUDADataFormats/EcalDigi/interface/DigisCollection.h

@@ -0,0 +1,24 @@
+#ifndef CUDADataFormats_EcalDigi_interface_DigisCollection_h
+#define CUDADataFormats_EcalDigi_interface_DigisCollection_h
+
+#include "CUDADataFormats/CaloCommon/interface/Common.h"
+
+namespace ecal {
+
+  template <typename StoragePolicy>
+  struct DigisCollection : public ::calo::common::AddSize<typename StoragePolicy::TagType> {
+    DigisCollection() = default;
+    DigisCollection(DigisCollection const &) = default;
+    DigisCollection &operator=(DigisCollection const &) = default;
+
+    DigisCollection(DigisCollection &&) = default;
+    DigisCollection &operator=(DigisCollection &&) = default;
+
+    // stride is statically known
+    typename StoragePolicy::template StorageSelector<uint32_t>::type ids;
+    typename StoragePolicy::template StorageSelector<uint16_t>::type data;
+  };
+
+}  // namespace ecal
+
+#endif  // CUDADataFormats_EcalDigi_interface_DigisCollection_h

CUDADataFormats/EcalDigi/src/classes.h

@@ -0,0 +1,3 @@
+#include "DataFormats/Common/interface/Wrapper.h"
+#include "CUDADataFormats/Common/interface/Product.h"
+#include "CUDADataFormats/EcalDigi/interface/DigisCollection.h"

CUDADataFormats/EcalDigi/src/classes_def.xml

@@ -0,0 +1,6 @@
+<lcgdict>
+    <class name="cms::cuda::Product<ecal::DigisCollection<calo::common::ViewStoragePolicy>>" persistent="false" />
+    <class name="cms::cuda::Product<ecal::DigisCollection<calo::common::DevStoragePolicy>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::DigisCollection<calo::common::ViewStoragePolicy>>>" persistent="false"/>
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::DigisCollection<calo::common::DevStoragePolicy>>>" persistent="false"/>
+</lcgdict>

CUDADataFormats/EcalRecHitSoA/BuildFile.xml

@@ -0,0 +1,10 @@
+<use name="cuda"/>
+<use name="CUDADataFormats/Common"/>
+<use name="CUDADataFormats/CaloCommon"/>
+<use name="DataFormats/Common"/>
+<use name="DataFormats/EcalDigi"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+  <lib   name="1"/>
+</export>