18 #include <type_traits>
20 #ifdef ALPAKA_ACC_SYCL_ENABLED
22 # include <sycl/sycl.hpp>
27 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
28 class BufGenericSycl :
public internal::ViewAccessOps<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
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 "
35 static_assert(!std::is_const_v<TIdx>,
"The idx type of the buffer can not be const!");
38 template<
typename TExtent,
typename Deleter>
39 BufGenericSycl(DevGenericSycl<TPlatform>
const& dev, TElem*
const pMem, Deleter deleter, TExtent
const& extent)
42 , m_spMem(pMem, std::move(deleter))
47 TDim::value == Dim<TExtent>::value,
48 "The dimensionality of TExtent and the dimensionality of the TDim template parameter have to be "
52 std::is_same_v<TIdx, Idx<TExtent>>,
53 "The idx type of TExtent and the TIdx template parameter have to be identical!");
56 DevGenericSycl<TPlatform> m_dev;
57 Vec<TDim, TIdx> m_extentElements;
58 std::shared_ptr<TElem> m_spMem;
65 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
66 struct DevType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
68 using type = DevGenericSycl<TPlatform>;
72 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
73 struct GetDev<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
75 static auto getDev(BufGenericSycl<TElem, TDim, TIdx, TPlatform>
const& buf)
82 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
83 struct DimType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
89 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
90 struct ElemType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
96 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
97 struct GetExtents<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
99 auto operator()(BufGenericSycl<TElem, TDim, TIdx, TPlatform>
const& buf)
const
101 return buf.m_extentElements;
106 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
107 struct GetPtrNative<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
109 static auto getPtrNative(BufGenericSycl<TElem, TDim, TIdx, TPlatform>
const& buf) -> TElem
const*
111 return buf.m_spMem.get();
114 static auto getPtrNative(BufGenericSycl<TElem, TDim, TIdx, TPlatform>& buf) -> TElem*
116 return buf.m_spMem.get();
121 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
122 struct GetPtrDev<BufGenericSycl<TElem, TDim, TIdx, TPlatform>, DevGenericSycl<TPlatform>>
125 BufGenericSycl<TElem, TDim, TIdx, TPlatform>
const& buf,
126 DevGenericSycl<TPlatform>
const& dev) -> TElem
const*
130 return buf.m_spMem.get();
134 throw std::runtime_error(
"The buffer is not accessible from the given device!");
138 static auto getPtrDev(BufGenericSycl<TElem, TDim, TIdx, TPlatform>& buf, DevGenericSycl<TPlatform>
const& dev)
143 return buf.m_spMem.get();
147 throw std::runtime_error(
"The buffer is not accessible from the given device!");
153 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
154 struct BufAlloc<TElem, TDim, TIdx, DevGenericSycl<TPlatform>>
156 template<
typename TExtent>
157 static auto allocBuf(DevGenericSycl<TPlatform>
const& dev, TExtent
const& extent)
158 -> BufGenericSycl<TElem, TDim, TIdx, TPlatform>
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)
167 auto const width =
getWidth(extent);
169 auto const widthBytes = width *
static_cast<TIdx
>(
sizeof(TElem));
170 std::cout << __func__ <<
" ew: " << width <<
" ewb: " << widthBytes <<
'\n';
172 else if constexpr(TDim::value == 2)
174 auto const width =
getWidth(extent);
177 auto const widthBytes = width *
static_cast<TIdx
>(
sizeof(TElem));
178 std::cout << __func__ <<
" ew: " << width <<
" eh: " << height <<
" ewb: " << widthBytes
179 <<
" pitch: " << widthBytes <<
'\n';
181 else if constexpr(TDim::value == 3)
183 auto const width =
getWidth(extent);
185 auto const depth =
getDepth(extent);
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';
193 auto const& [nativeDev, nativeContext] = dev.getNativeHandle();
194 TElem* memPtr = sycl::malloc_device<TElem>(
198 auto deleter = [ctx = nativeContext](TElem* ptr) {
sycl::free(ptr, ctx); };
200 return BufGenericSycl<TElem, TDim, TIdx, TPlatform>(dev, memPtr, std::move(deleter), extent);
205 template<
typename TDim,
typename TPlatform>
206 struct HasAsyncBufSupport<TDim, DevGenericSycl<TPlatform>> : std::false_type
211 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
212 struct GetOffsets<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
214 auto operator()(BufGenericSycl<TElem, TDim, TIdx, TPlatform>
const&)
const -> Vec<TDim, TIdx>
221 template<
typename TPlatform,
typename TElem,
typename TDim,
typename TIdx>
222 struct BufAllocMapped
224 template<
typename TExtent>
225 static auto allocMappedBuf(DevCpu
const& host, TPlatform
const& platform, TExtent
const& extent)
226 -> BufCpu<TElem, TDim, TIdx>
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); };
236 return BufCpu<TElem, TDim, TIdx>(host, memPtr, std::move(deleter), extent);
241 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
242 struct IdxType<BufGenericSycl<TElem, TDim, TIdx, TPlatform>>
248 template<
typename TElem,
typename TDim,
typename TIdx,
typename TPlatform>
249 struct GetPtrDev<BufCpu<TElem, TDim, TIdx>, DevGenericSycl<TPlatform>>
251 static auto getPtrDev(BufCpu<TElem, TDim, TIdx>
const& buf, DevGenericSycl<TPlatform>
const&) -> TElem
const*
256 static auto getPtrDev(BufCpu<TElem, TDim, TIdx>& buf, DevGenericSycl<TPlatform>
const&) -> TElem*
#define ALPAKA_DEBUG_MINIMAL_LOG_SCOPE
ALPAKA_NO_HOST_ACC_WARNING static constexpr ALPAKA_FN_HOST_ACC auto zeros() -> Vec< TDim, TVal >
Zero value constructor.
The alpaka accelerator library.
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getExtentProduct(T const &object) -> Idx< T >
ALPAKA_FN_HOST auto free(TAlloc const &alloc, T const *const ptr) -> void
Frees the memory identified by the given pointer.
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getHeight(TExtent const &extent=TExtent()) -> Idx< TExtent >
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.
ALPAKA_FN_HOST auto getPtrNative(TView const &view) -> Elem< TView > const *
Gets the native pointer of the memory view.
ALPAKA_FN_HOST auto allocBuf(TDev const &dev, TExtent const &extent=TExtent())
Allocates memory on the given device.
ALPAKA_FN_HOST auto getDev(T const &t)
ALPAKA_FN_HOST auto getPtrDev(TView const &view, TDev const &dev) -> Elem< TView > const *
Gets the pointer to the view on the given device.
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getDepth(TExtent const &extent=TExtent()) -> Idx< TExtent >
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getWidth(TExtent const &extent=TExtent()) -> Idx< TExtent >
ALPAKA_NO_HOST_ACC_WARNING constexpr ALPAKA_FN_HOST_ACC auto getExtentVecEnd(T const &object={}) -> Vec< TDim, Idx< T >>