Synchronise with CMSSW_11_1_0

CUDADataFormats/BeamSpot/BuildFile.xml

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

CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h

@@ -0,0 +1,32 @@
+#ifndef CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h
+#define CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+
+#include <cuda_runtime.h>
+
+class BeamSpotCUDA {
+public:
+  // alignas(128) doesn't really make sense as there is only one
+  // beamspot per event?
+  struct Data {
+    float x, y, z;  // position
+    // TODO: add covariance matrix
+
+    float sigmaZ;
+    float beamWidthX, beamWidthY;
+    float dxdz, dydz;
+    float emittanceX, emittanceY;
+    float betaStar;
+  };
+
+  BeamSpotCUDA() = default;
+  BeamSpotCUDA(Data const* data_h, cudaStream_t stream);
+
+  Data const* data() const { return data_d_.get(); }
+
+private:
+  cms::cuda::device::unique_ptr<Data> data_d_;
+};
+
+#endif

CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc

@@ -0,0 +1,9 @@
+#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h"
+
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+
+BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cudaStream_t stream) {
+  data_d_ = cms::cuda::make_device_unique<Data>(stream);
+  cudaCheck(cudaMemcpyAsync(data_d_.get(), data_h, sizeof(Data), cudaMemcpyHostToDevice, stream));
+}

CUDADataFormats/BeamSpot/src/classes.h

@@ -0,0 +1,8 @@
+#ifndef CUDADataFormats_BeamSpot_classes_h
+#define CUDADataFormats_BeamSpot_classes_h
+
+#include "CUDADataFormats/Common/interface/Product.h"
+#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif

CUDADataFormats/BeamSpot/src/classes_def.xml

@@ -0,0 +1,4 @@
+<lcgdict>
+  <class name="cms::cuda::Product<BeamSpotCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<BeamSpotCUDA>>" persistent="false"/>
+</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,7 @@
+<use name="CUDADataFormats/Common"/>
+<use name="DataFormats/Common"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+  <lib   name="1"/>
+</export>

CUDADataFormats/EcalDigi/interface/DigisCollection.h

@@ -0,0 +1,27 @@
+#ifndef CUDADataFormats_EcalDigi_interface_DigisCollection_h
+#define CUDADataFormats_EcalDigi_interface_DigisCollection_h
+
+namespace ecal {
+
+  //
+  // this is basically a view
+  // it does not own the actual memory -> does not reclaim
+  //
+  struct DigisCollection {
+    DigisCollection() = default;
+    DigisCollection(uint32_t *ids, uint16_t *data, uint32_t ndigis) : ids{ids}, data{data}, ndigis{ndigis} {}
+    DigisCollection(DigisCollection const &) = default;
+    DigisCollection &operator=(DigisCollection const &) = default;
+
+    DigisCollection(DigisCollection &&) = default;
+    DigisCollection &operator=(DigisCollection &&) = default;
+
+    // stride is statically known
+    uint32_t *ids = nullptr;
+    uint16_t *data = nullptr;
+    uint32_t ndigis;
+  };
+
+}  // 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,4 @@
+<lcgdict>
+    <class name="cms::cuda::Product<ecal::DigisCollection>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::DigisCollection>>" persistent="false"/>
+</lcgdict>

CUDADataFormats/EcalRecHitSoA/BuildFile.xml

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