CMS 3D CMS Logo

host_unique_ptr.h
Go to the documentation of this file.
1 #ifndef HeterogeneousCore_CUDAUtilities_interface_host_unique_ptr_h
2 #define HeterogeneousCore_CUDAUtilities_interface_host_unique_ptr_h
3 
4 #include <cstdlib>
5 #include <memory>
6 #include <functional>
7 
9 
10 namespace cms {
11  namespace cuda {
12  namespace host {
13  namespace impl {
14 
15  enum class MemoryType : bool {
16  kDefault = false,
17  kPinned = true,
18  };
19 
20  // Custom deleter for host memory, with an internal state to distinguish pageable and pinned host memory
21  class HostDeleter {
22  public:
23  // The default constructor is needed by the default constructor of unique_ptr<T, HostDeleter>,
24  // which is needed by the default constructor of HostProduct<T>, which is needed by the ROOT dictionary
27 
28  void operator()(void *ptr) {
29  if (type_ == MemoryType::kPinned) {
31  } else {
32  std::free(ptr);
33  }
34  }
35 
36  private:
38  };
39 
40  } // namespace impl
41 
42  template <typename T>
43  using unique_ptr = std::unique_ptr<T, impl::HostDeleter>;
44 
45  namespace impl {
46  template <typename T>
49  };
50  template <typename T>
53  };
54  template <typename T, size_t N>
56  struct bounded_array {};
57  };
58  } // namespace impl
59  } // namespace host
60 
61  // Allocate pageable host memory
62  template <typename T>
65  "Allocating with non-trivial constructor on the host memory is not supported");
66  // Allocate a buffer aligned to 128 bytes, to match the CUDA cache line size
67  const size_t alignment = 128;
68  // std::aligned_alloc() requires the size to be a multiple of the alignment
69  const size_t size = (sizeof(T) + alignment - 1) / alignment * alignment;
70  void *mem = std::aligned_alloc(alignment, size);
71  return typename host::impl::make_host_unique_selector<T>::non_array{reinterpret_cast<T *>(mem),
73  }
74 
75  template <typename T>
77  using element_type = typename std::remove_extent<T>::type;
79  "Allocating with non-trivial constructor on the host memory is not supported");
80  // Allocate a buffer aligned to 128 bytes, to match the CUDA cache line size
81  const size_t alignment = 128;
82  // std::aligned_alloc() requires the size to be a multiple of the alignment
83  const size_t size = (n * sizeof(element_type) + alignment - 1) / alignment * alignment;
84  void *mem = std::aligned_alloc(alignment, size);
85  return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem),
87  }
88 
89  // Allocate pinned host memory
90  template <typename T>
93  "Allocating with non-trivial constructor on the host memory is not supported");
94  void *mem = allocate_host(sizeof(T), stream);
95  return typename host::impl::make_host_unique_selector<T>::non_array{reinterpret_cast<T *>(mem), //
97  }
98 
99  template <typename T>
101  using element_type = typename std::remove_extent<T>::type;
103  "Allocating with non-trivial constructor on the host memory is not supported");
104  void *mem = allocate_host(n * sizeof(element_type), stream);
105  return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem),
107  }
108 
109  // Arrays of known bounds are not supported by std::unique_ptr
110  template <typename T, typename... Args>
111  typename host::impl::make_host_unique_selector<T>::bounded_array make_host_unique(Args &&...) = delete;
112 
113  // No check for the trivial constructor, make it clear in the interface
114  template <typename T>
116  void *mem = allocate_host(sizeof(T), stream);
117  return typename host::impl::make_host_unique_selector<T>::non_array{reinterpret_cast<T *>(mem), //
119  }
120 
121  template <typename T>
123  size_t n, cudaStream_t stream) {
124  using element_type = typename std::remove_extent<T>::type;
125  void *mem = allocate_host(n * sizeof(element_type), stream);
126  return typename host::impl::make_host_unique_selector<T>::unbounded_array{reinterpret_cast<element_type *>(mem),
128  }
129 
130  // Arrays of known bounds are not supported by std::unique_ptr
131  template <typename T, typename... Args>
132  typename host::impl::make_host_unique_selector<T>::bounded_array make_host_unique_uninitialized(Args &&...) = delete;
133 
134  } // namespace cuda
135 } // namespace cms
136 
137 #endif
string host
Definition: query.py:115
host::impl::make_host_unique_selector< T >::non_array make_host_unique()
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
void free_host(void *ptr)
cms::cuda::host::unique_ptr< T > non_array
host::impl::make_host_unique_selector< T >::non_array make_host_unique_uninitialized(cudaStream_t stream)
Namespace of DDCMS conversion namespace.
#define N
Definition: blowfish.cc:9
cms::cuda::host::unique_ptr< T[]> unbounded_array
__host__ __device__ VT uint32_t size
Definition: prefixScan.h:47
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int n
std::unique_ptr< T, impl::HostDeleter > unique_ptr
void * allocate_host(size_t nbytes, cudaStream_t stream)
long double T