[WIP] Follow-ups to the CUDA framework PR review

CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h

@@ -26,7 +26,7 @@ class BeamSpotCUDA {
   Data const* data() const { return data_d_.get(); }
 
 private:
-  cudautils::device::unique_ptr<Data> data_d_;
+  cms::cuda::device::unique_ptr<Data> data_d_;
 };
 
 #endif

CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc

@@ -4,6 +4,6 @@
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 
 BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cudaStream_t stream) {
-  data_d_ = cudautils::make_device_unique<Data>(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

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

CUDADataFormats/BeamSpot/src/classes_def.xml

@@ -1,4 +1,4 @@
 <lcgdict>
-  <class name="CUDAProduct<BeamSpotCUDA>" persistent="false"/>
-  <class name="edm::Wrapper<CUDAProduct<BeamSpotCUDA>>" persistent="false"/>
+  <class name="cms::cuda::Product<BeamSpotCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<BeamSpotCUDA>>" persistent="false"/>
 </lcgdict>

CUDADataFormats/Common/BuildFile.xml

@@ -1,5 +1,7 @@
+<iftool name="cuda-gcc-support">
 <use name="HeterogeneousCore/CUDAUtilities"/>
 
 <export>
     <lib name="1"/>
 </export>
+</iftool>

CUDADataFormats/Common/interface/HeterogeneousSoA.h

@@ -19,8 +19,8 @@ class HeterogeneousSoA {
   HeterogeneousSoA(HeterogeneousSoA &&) = default;
   HeterogeneousSoA &operator=(HeterogeneousSoA &&) = default;
 
-  explicit HeterogeneousSoA(cudautils::device::unique_ptr<T> &&p) : dm_ptr(std::move(p)) {}
-  explicit HeterogeneousSoA(cudautils::host::unique_ptr<T> &&p) : hm_ptr(std::move(p)) {}
+  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()); }
@@ -36,74 +36,74 @@ class HeterogeneousSoA {
   auto *operator-> () { return get(); }
 
   // in reality valid only for GPU version...
-  cudautils::host::unique_ptr<T> toHostAsync(cudaStream_t stream) const {
+  cms::cuda::host::unique_ptr<T> toHostAsync(cudaStream_t stream) const {
     assert(dm_ptr);
-    auto ret = cudautils::make_host_unique<T>(stream);
+    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
-  cudautils::device::unique_ptr<T> dm_ptr;  //!
-  cudautils::host::unique_ptr<T> hm_ptr;    //!
+  cms::cuda::device::unique_ptr<T> dm_ptr;  //!
+  cms::cuda::host::unique_ptr<T> hm_ptr;    //!
   std::unique_ptr<T> std_ptr;               //!
 };
 
 namespace cudaCompat {
 
   struct GPUTraits {
     template <typename T>
-    using unique_ptr = cudautils::device::unique_ptr<T>;
+    using unique_ptr = cms::cuda::device::unique_ptr<T>;
 
     template <typename T>
     static auto make_unique(cudaStream_t stream) {
-      return cudautils::make_device_unique<T>(stream);
+      return cms::cuda::make_device_unique<T>(stream);
     }
 
     template <typename T>
     static auto make_unique(size_t size, cudaStream_t stream) {
-      return cudautils::make_device_unique<T>(size, stream);
+      return cms::cuda::make_device_unique<T>(size, stream);
     }
 
     template <typename T>
     static auto make_host_unique(cudaStream_t stream) {
-      return cudautils::make_host_unique<T>(stream);
+      return cms::cuda::make_host_unique<T>(stream);
     }
 
     template <typename T>
     static auto make_device_unique(cudaStream_t stream) {
-      return cudautils::make_device_unique<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 cudautils::make_device_unique<T>(size, stream);
+      return cms::cuda::make_device_unique<T>(size, stream);
     }
   };
 
   struct HostTraits {
     template <typename T>
-    using unique_ptr = cudautils::host::unique_ptr<T>;
+    using unique_ptr = cms::cuda::host::unique_ptr<T>;
 
     template <typename T>
     static auto make_unique(cudaStream_t stream) {
-      return cudautils::make_host_unique<T>(stream);
+      return cms::cuda::make_host_unique<T>(stream);
     }
 
     template <typename T>
     static auto make_host_unique(cudaStream_t stream) {
-      return cudautils::make_host_unique<T>(stream);
+      return cms::cuda::make_host_unique<T>(stream);
     }
 
     template <typename T>
     static auto make_device_unique(cudaStream_t stream) {
-      return cudautils::make_device_unique<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 cudautils::make_device_unique<T>(size, stream);
+      return cms::cuda::make_device_unique<T>(size, stream);
     }
   };
 
@@ -158,7 +158,7 @@ class HeterogeneousSoAImpl {
 
   T *get() { return m_ptr.get(); }
 
-  cudautils::host::unique_ptr<T> toHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<T> toHostAsync(cudaStream_t stream) const;
 
 private:
   unique_ptr<T> m_ptr;  //!
@@ -171,8 +171,8 @@ HeterogeneousSoAImpl<T, Traits>::HeterogeneousSoAImpl(cudaStream_t stream) {
 
 // in reality valid only for GPU version...
 template <typename T, typename Traits>
-cudautils::host::unique_ptr<T> HeterogeneousSoAImpl<T, Traits>::toHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<T>(stream);
+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;
 }

CUDADataFormats/Common/interface/HostProduct.h

@@ -12,7 +12,7 @@ class HostProduct {
   HostProduct(HostProduct&&) = default;
   HostProduct& operator=(HostProduct&&) = default;
 
-  explicit HostProduct(cudautils::host::unique_ptr<T>&& p) : hm_ptr(std::move(p)) {}
+  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(); }
@@ -22,7 +22,7 @@ class HostProduct {
   auto const* operator-> () const { return get(); }
 
 private:
-  cudautils::host::unique_ptr<T> hm_ptr;  //!
+  cms::cuda::host::unique_ptr<T> hm_ptr;  //!
   std::unique_ptr<T> std_ptr;             //!
 };
 

CUDADataFormats/Common/interface/Product.h

@@ -0,0 +1,60 @@
+#ifndef CUDADataFormats_Common_Product_h
+#define CUDADataFormats_Common_Product_h
+
+#include <memory>
+
+#include "CUDADataFormats/Common/interface/ProductBase.h"
+
+namespace edm {
+  template <typename T>
+  class Wrapper;
+}
+
+namespace cms {
+  namespace cuda {
+    namespace impl {
+      class ScopedContextGetterBase;
+    }
+
+    /**
+     * The purpose of this class is to wrap CUDA data to edm::Event in a
+     * way which forces correct use of various utilities.
+     *
+     * The non-default construction has to be done with cms::cuda::ScopedContext
+     * (in order to properly register the CUDA event).
+     *
+     * The default constructor is needed only for the ROOT dictionary generation.
+     *
+     * The CUDA event is in practice needed only for stream-stream
+     * synchronization, but someone with long-enough lifetime has to own
+     * it. Here is a somewhat natural place. If overhead is too much, we
+     * can use them only where synchronization between streams is needed.
+     */
+    template <typename T>
+    class Product : public ProductBase {
+    public:
+      Product() = default;  // Needed only for ROOT dictionary generation
+
+      Product(const Product&) = delete;
+      Product& operator=(const Product&) = delete;
+      Product(Product&&) = default;
+      Product& operator=(Product&&) = default;
+
+    private:
+      friend class impl::ScopedContextGetterBase;
+      friend class ScopedContextProduce;
+      friend class edm::Wrapper<Product<T>>;
+
+      explicit Product(int device, SharedStreamPtr stream, SharedEventPtr event, T data)
+          : ProductBase(device, std::move(stream), std::move(event)), data_(std::move(data)) {}
+
+      template <typename... Args>
+      explicit Product(int device, SharedStreamPtr stream, SharedEventPtr event, Args&&... args)
+          : ProductBase(device, std::move(stream), std::move(event)), data_(std::forward<Args>(args)...) {}
+
+      T data_;  //!
+    };
+  }  // namespace cuda
+}  // namespace cms
+
+#endif