CMS 3D CMS Logo

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
Likely.h
allocate_device.h
cms::cuda::make_device_unique
device::impl::make_device_unique_selector< T >::non_array make_device_unique(cudaStream_t stream)
Definition: device_unique_ptr.h:52
cms::cuda::stream
cudaStream_t stream
Definition: HistoContainer.h:57
mem
uint16_t mem[nChs][nEvts]
Definition: recycleTccEmu.cc:13
cms::cuda::device::impl::DeviceDeleter::DeviceDeleter
DeviceDeleter()=default
N
#define N
Definition: blowfish.cc:9
cms::cuda::currentDevice
int currentDevice()
Definition: currentDevice.h:10
cms::cuda::device::impl::DeviceDeleter::operator()
void operator()(void *ptr)
Definition: device_unique_ptr.h:21
cms::cuda::n
cudaStream_t T uint32_t const T *__restrict__ const uint32_t *__restrict__ uint32_t int cudaStream_t Func __host__ __device__ V int n
Definition: HistoContainer.h:124
gainCalibHelper::gainCalibPI::type
type
Definition: SiPixelGainCalibHelper.h:40
cms::cuda::device::impl::make_device_unique_selector::non_array
cms::cuda::device::unique_ptr< T > non_array
Definition: device_unique_ptr.h:38
cms::cuda::device::impl::DeviceDeleter::DeviceDeleter
DeviceDeleter(int device)
Definition: device_unique_ptr.h:19
cms::cuda::free_device
void free_device(int device, void *ptr)
Definition: allocate_device.cc:31
cms::cuda::make_device_unique_uninitialized
device::impl::make_device_unique_selector< T >::non_array make_device_unique_uninitialized(cudaStream_t stream)
Definition: device_unique_ptr.h:78
cms::cuda::device::impl::make_device_unique_selector
Definition: device_unique_ptr.h:37
cms::cuda::device::unique_ptr
std::unique_ptr< T, impl::DeviceDeleter > unique_ptr
Definition: device_unique_ptr.h:33
impl
Definition: trackAlgoPriorityOrder.h:18
T
long double T
Definition: Basic3DVectorLD.h:48
LIKELY
#define LIKELY(x)
Definition: Likely.h:20
relativeConstraints.value
value
Definition: relativeConstraints.py:53
cms::cuda::device::impl::make_device_unique_selector< T[]>::unbounded_array
cms::cuda::device::unique_ptr< T[]> unbounded_array
Definition: device_unique_ptr.h:42
currentDevice.h
ecalDigis_cff.cuda
cuda
Definition: ecalDigis_cff.py:35
cms::cuda::allocate_device
void * allocate_device(int dev, size_t nbytes, cudaStream_t stream)
Definition: allocate_device.cc:16
cms::cuda::device::impl::DeviceDeleter
Definition: device_unique_ptr.h:16
cms::cuda::device::impl::DeviceDeleter::device_
int device_
Definition: device_unique_ptr.h:28
cms
Namespace of DDCMS conversion namespace.
Definition: ProducerAnalyzer.cc:21