CMS 3D CMS Logo

 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Properties Friends Macros Groups Pages
device_unique_ptr.h
Go to the documentation of this file.
1 #ifndef HeterogeneousCore_CUDAUtilities_interface_device_unique_ptr_h
2 #define HeterogeneousCore_CUDAUtilities_interface_device_unique_ptr_h
3 
4 #include <memory>
5 #include <functional>
6 
10 
11 namespace cms {
12  namespace cuda {
13  namespace device {
14  namespace impl {
15  // Additional layer of types to distinguish from host::unique_ptr
16  class DeviceDeleter {
17  public:
18  DeviceDeleter() = default; // for edm::Wrapper
19  DeviceDeleter(int device) : device_{device} {}
20 
21  void operator()(void *ptr) {
22  if (LIKELY(device_ >= 0)) {
23  free_device(device_, ptr);
24  }
25  }
26 
27  private:
28  int device_ = -1;
29  };
30  } // namespace impl
31 
32  template <typename T>
33  using unique_ptr = std::unique_ptr<T, impl::DeviceDeleter>;
34 
35  namespace impl {
36  template <typename T>
39  };
40  template <typename T>
43  };
44  template <typename T, size_t N>
46  struct bounded_array {};
47  };
48  } // namespace impl
49  } // namespace device
50 
51  template <typename T>
54  "Allocating with non-trivial constructor on the device memory is not supported");
55  int dev = currentDevice();
56  void *mem = allocate_device(dev, sizeof(T), stream);
57  return typename device::impl::make_device_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),
59  }
60 
61  template <typename T>
63  cudaStream_t stream) {
64  using element_type = typename std::remove_extent<T>::type;
66  "Allocating with non-trivial constructor on the device memory is not supported");
67  int dev = currentDevice();
68  void *mem = allocate_device(dev, n * sizeof(element_type), stream);
70  reinterpret_cast<element_type *>(mem), device::impl::DeviceDeleter{dev}};
71  }
72 
73  template <typename T, typename... Args>
74  typename device::impl::make_device_unique_selector<T>::bounded_array make_device_unique(Args &&...) = delete;
75 
76  // No check for the trivial constructor, make it clear in the interface
77  template <typename T>
79  cudaStream_t stream) {
80  int dev = currentDevice();
81  void *mem = allocate_device(dev, sizeof(T), stream);
82  return typename device::impl::make_device_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),
84  }
85 
86  template <typename T>
88  size_t n, cudaStream_t stream) {
89  using element_type = typename std::remove_extent<T>::type;
90  int dev = currentDevice();
91  void *mem = allocate_device(dev, n * sizeof(element_type), stream);
93  reinterpret_cast<element_type *>(mem), device::impl::DeviceDeleter{dev}};
94  }
95 
96  template <typename T, typename... Args>
97  typename device::impl::make_device_unique_selector<T>::bounded_array make_device_unique_uninitialized(Args &&...) =
98  delete;
99  } // namespace cuda
100 } // namespace cms
101 
102 #endif
#define LIKELY(x)
Definition: Likely.h:20
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type * mem
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t stream
std::unique_ptr< T, impl::DeviceDeleter > unique_ptr
device::impl::make_device_unique_selector< T >::non_array make_device_unique_uninitialized(cudaStream_t stream)
void free_device(int device, void *ptr)
#define N
Definition: blowfish.cc:9
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int n
void * allocate_device(int dev, size_t nbytes, cudaStream_t stream)
device::impl::make_device_unique_selector< T >::non_array make_device_unique(cudaStream_t stream)
long double T
int currentDevice()
Definition: currentDevice.h:10