alpaka
Abstraction Library for Parallel Kernel Acceleration
BufGenericSycl.hpp
Go to the documentation of this file.
1 /* Copyright 2023 Jan Stephan, Luca Ferragina, Aurora Perego, Andrea Bocci
2  * SPDX-License-Identifier: MPL-2.0
3  */
4 
5 #pragma once
6 
7 #include "alpaka/core/Sycl.hpp"
9 #include "alpaka/dev/Traits.hpp"
11 #include "alpaka/dim/Traits.hpp"
15 #include "alpaka/vec/Vec.hpp"
16 
17 #include <memory>
18 #include <type_traits>
19 
20 #ifdef ALPAKA_ACC_SYCL_ENABLED
21 
22 # include <sycl/sycl.hpp>
23 
24 namespace alpaka
25 {
26  //! The SYCL memory buffer.
27  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
28  class BufGenericSycl : public internal::ViewAccessOps<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
29  {
30  public:
31  static_assert(
32  !std::is_const_v<TElem>,
33  "The elem type of the buffer can not be const because the C++ Standard forbids containers of const "
34  "elements!");
35  static_assert(!std::is_const_v<TIdx>, "The idx type of the buffer can not be const!");
36 
37  //! Constructor
38  template<typename TExtent, typename Deleter>
39  BufGenericSycl(DevGenericSycl<TPlatform> const& dev, TElem* const pMem, Deleter deleter, TExtent const& extent)
40  : m_dev{dev}
41  , m_extentElements{getExtentVecEnd<TDim>(extent)}
42  , m_spMem(pMem, std::move(deleter))
43  {
45 
46  static_assert(
47  TDim::value == Dim<TExtent>::value,
48  "The dimensionality of TExtent and the dimensionality of the TDim template parameter have to be "
49  "identical!");
50 
51  static_assert(
52  std::is_same_v<TIdx, Idx<TExtent>>,
53  "The idx type of TExtent and the TIdx template parameter have to be identical!");
54  }
55 
56  DevGenericSycl<TPlatform> m_dev;
57  Vec<TDim, TIdx> m_extentElements;
58  std::shared_ptr<TElem> m_spMem;
59  };
60 } // namespace alpaka
61 
62 namespace alpaka::trait
63 {
64  //! The BufGenericSycl device type trait specialization.
65  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
66  struct DevType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
67  {
68  using type = DevGenericSycl<TPlatform>;
69  };
70 
71  //! The BufGenericSycl device get trait specialization.
72  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
73  struct GetDev<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
74  {
75  static auto getDev(BufGenericSycl<TElem, TDim, TIdx, TPlatform> const& buf)
76  {
77  return buf.m_dev;
78  }
79  };
80 
81  //! The BufGenericSycl dimension getter trait specialization.
82  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
83  struct DimType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
84  {
85  using type = TDim;
86  };
87 
88  //! The BufGenericSycl memory element type get trait specialization.
89  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
90  struct ElemType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
91  {
92  using type = TElem;
93  };
94 
95  //! The BufGenericSycl extent get trait specialization.
96  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
97  struct GetExtents<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
98  {
99  auto operator()(BufGenericSycl<TElem, TDim, TIdx, TPlatform> const& buf) const
100  {
101  return buf.m_extentElements;
102  }
103  };
104 
105  //! The BufGenericSycl native pointer get trait specialization.
106  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
107  struct GetPtrNative<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
108  {
109  static auto getPtrNative(BufGenericSycl<TElem, TDim, TIdx, TPlatform> const& buf) -> TElem const*
110  {
111  return buf.m_spMem.get();
112  }
113 
114  static auto getPtrNative(BufGenericSycl<TElem, TDim, TIdx, TPlatform>& buf) -> TElem*
115  {
116  return buf.m_spMem.get();
117  }
118  };
119 
120  //! The BufGenericSycl pointer on device get trait specialization.
121  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
122  struct GetPtrDev<BufGenericSycl<TElem, TDim, TIdx, TPlatform>, DevGenericSycl<TPlatform>>
123  {
124  static auto getPtrDev(
125  BufGenericSycl<TElem, TDim, TIdx, TPlatform> const& buf,
126  DevGenericSycl<TPlatform> const& dev) -> TElem const*
127  {
128  if(dev == getDev(buf))
129  {
130  return buf.m_spMem.get();
131  }
132  else
133  {
134  throw std::runtime_error("The buffer is not accessible from the given device!");
135  }
136  }
137 
138  static auto getPtrDev(BufGenericSycl<TElem, TDim, TIdx, TPlatform>& buf, DevGenericSycl<TPlatform> const& dev)
139  -> TElem*
140  {
141  if(dev == getDev(buf))
142  {
143  return buf.m_spMem.get();
144  }
145  else
146  {
147  throw std::runtime_error("The buffer is not accessible from the given device!");
148  }
149  }
150  };
151 
152  //! The SYCL memory allocation trait specialization.
153  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
154  struct BufAlloc<TElem, TDim, TIdx, DevGenericSycl<TPlatform>>
155  {
156  template<typename TExtent>
157  static auto allocBuf(DevGenericSycl<TPlatform> const& dev, TExtent const& extent)
158  -> BufGenericSycl<TElem, TDim, TIdx, TPlatform>
159  {
161 
162 # if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
163  if constexpr(TDim::value == 0)
164  std::cout << __func__ << " ewb: " << sizeof(TElem) << '\n';
165  else if constexpr(TDim::value == 1)
166  {
167  auto const width = getWidth(extent);
168 
169  auto const widthBytes = width * static_cast<TIdx>(sizeof(TElem));
170  std::cout << __func__ << " ew: " << width << " ewb: " << widthBytes << '\n';
171  }
172  else if constexpr(TDim::value == 2)
173  {
174  auto const width = getWidth(extent);
175  auto const height = getHeight(extent);
176 
177  auto const widthBytes = width * static_cast<TIdx>(sizeof(TElem));
178  std::cout << __func__ << " ew: " << width << " eh: " << height << " ewb: " << widthBytes
179  << " pitch: " << widthBytes << '\n';
180  }
181  else if constexpr(TDim::value == 3)
182  {
183  auto const width = getWidth(extent);
184  auto const height = getHeight(extent);
185  auto const depth = getDepth(extent);
186 
187  auto const widthBytes = width * static_cast<TIdx>(sizeof(TElem));
188  std::cout << __func__ << " ew: " << width << " eh: " << height << " ed: " << depth
189  << " ewb: " << widthBytes << " pitch: " << widthBytes << '\n';
190  }
191 # endif
192 
193  auto const& [nativeDev, nativeContext] = dev.getNativeHandle();
194  TElem* memPtr = sycl::malloc_device<TElem>(
195  static_cast<std::size_t>(getExtentProduct(extent)),
196  nativeDev,
197  nativeContext);
198  auto deleter = [ctx = nativeContext](TElem* ptr) { sycl::free(ptr, ctx); };
199 
200  return BufGenericSycl<TElem, TDim, TIdx, TPlatform>(dev, memPtr, std::move(deleter), extent);
201  }
202  };
203 
204  //! The BufGenericSycl stream-ordered memory allocation capability trait specialization.
205  template<typename TDim, typename TPlatform>
206  struct HasAsyncBufSupport<TDim, DevGenericSycl<TPlatform>> : std::false_type
207  {
208  };
209 
210  //! The BufGenericSycl offset get trait specialization.
211  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
212  struct GetOffsets<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
213  {
214  auto operator()(BufGenericSycl<TElem, TDim, TIdx, TPlatform> const&) const -> Vec<TDim, TIdx>
215  {
216  return Vec<TDim, TIdx>::zeros();
217  }
218  };
219 
220  //! The pinned/mapped memory allocation trait specialization for the SYCL devices.
221  template<typename TPlatform, typename TElem, typename TDim, typename TIdx>
222  struct BufAllocMapped
223  {
224  template<typename TExtent>
225  static auto allocMappedBuf(DevCpu const& host, TPlatform const& platform, TExtent const& extent)
226  -> BufCpu<TElem, TDim, TIdx>
227  {
229 
230  // Allocate SYCL page-locked memory on the host, mapped into the TPlatform address space and
231  // accessible to all devices in the TPlatform.
232  auto ctx = platform.syclContext();
233  TElem* memPtr = sycl::malloc_host<TElem>(static_cast<std::size_t>(getExtentProduct(extent)), ctx);
234  auto deleter = [ctx](TElem* ptr) { sycl::free(ptr, ctx); };
235 
236  return BufCpu<TElem, TDim, TIdx>(host, memPtr, std::move(deleter), extent);
237  }
238  };
239 
240  //! The BufGenericSycl idx type trait specialization.
241  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
242  struct IdxType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
243  {
244  using type = TIdx;
245  };
246 
247  //! The BufCpu pointer on SYCL device get trait specialization.
248  template<typename TElem, typename TDim, typename TIdx, typename TPlatform>
249  struct GetPtrDev<BufCpu<TElem, TDim, TIdx>, DevGenericSycl<TPlatform>>
250  {
251  static auto getPtrDev(BufCpu<TElem, TDim, TIdx> const& buf, DevGenericSycl<TPlatform> const&) -> TElem const*
252  {
253  return getPtrNative(buf);
254  }
255 
256  static auto getPtrDev(BufCpu<TElem, TDim, TIdx>& buf, DevGenericSycl<TPlatform> const&) -> TElem*
257  {
258  return getPtrNative(buf);
259  }
260  };
261 } // namespace alpaka::trait
262 
264 # include "alpaka/mem/buf/sycl/Set.hpp"
265 
266 #endif
#define ALPAKA_DEBUG_MINIMAL_LOG_SCOPE
Definition: Debug.hpp:55
ALPAKA_NO_HOST_ACC_WARNING static constexpr ALPAKA_FN_HOST_ACC auto zeros() -> Vec< TDim, TVal >
Zero value constructor.
Definition: Vec.hpp:126
The accelerator traits.
The alpaka accelerator library.
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getExtentProduct(T const &object) -> Idx< T >
Definition: Traits.hpp:134
ALPAKA_FN_HOST auto free(TAlloc const &alloc, T const *const ptr) -> void
Frees the memory identified by the given pointer.
Definition: Traits.hpp:41
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getHeight(TExtent const &extent=TExtent()) -> Idx< TExtent >
Definition: Traits.hpp:108
ALPAKA_FN_HOST auto allocMappedBuf(DevCpu const &host, TPlatform const &platform, TExtent const &extent=TExtent())
Allocates pinned/mapped host memory, accessible by all devices in the given platform.
Definition: Traits.hpp:138
ALPAKA_FN_HOST auto getPtrNative(TView const &view) -> Elem< TView > const *
Gets the native pointer of the memory view.
Definition: Traits.hpp:139
ALPAKA_FN_HOST auto allocBuf(TDev const &dev, TExtent const &extent=TExtent())
Allocates memory on the given device.
Definition: Traits.hpp:64
ALPAKA_FN_HOST auto getDev(T const &t)
Definition: Traits.hpp:68
ALPAKA_FN_HOST auto getPtrDev(TView const &view, TDev const &dev) -> Elem< TView > const *
Gets the pointer to the view on the given device.
Definition: Traits.hpp:160
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getDepth(TExtent const &extent=TExtent()) -> Idx< TExtent >
Definition: Traits.hpp:121
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getWidth(TExtent const &extent=TExtent()) -> Idx< TExtent >
Definition: Traits.hpp:95
ALPAKA_NO_HOST_ACC_WARNING constexpr ALPAKA_FN_HOST_ACC auto getExtentVecEnd(T const &object={}) -> Vec< TDim, Idx< T >>
Definition: Traits.hpp:78