38 #include <type_traits>
40 #ifdef ALPAKA_ACC_SYCL_ENABLED
42 # include <sycl/sycl.hpp>
49 template<
typename TDim,
typename TIdx>
51 :
public WorkDivGenericSycl<TDim, TIdx>
52 ,
public gb::IdxGbGenericSycl<TDim, TIdx>
53 ,
public bt::IdxBtGenericSycl<TDim, TIdx>
54 ,
public AtomicHierarchy<AtomicGenericSycl, AtomicGenericSycl, AtomicGenericSycl>
55 ,
public math::MathGenericSycl
56 ,
public BlockSharedMemDynGenericSycl
57 ,
public BlockSharedMemStGenericSycl
58 ,
public BlockSyncGenericSycl<TDim>
59 ,
public IntrinsicGenericSycl
60 ,
public MemFenceGenericSycl
61 # ifdef ALPAKA_DISABLE_VENDOR_RNG
62 ,
public rand::RandDefault
64 ,
public rand::RandGenericSycl<TDim>
66 ,
public warp::WarpGenericSycl<TDim>
68 static_assert(TDim::value > 0,
"The SYCL accelerator must have a dimension greater than zero.");
71 AccGenericSycl(AccGenericSycl
const&) =
delete;
72 AccGenericSycl(AccGenericSycl&&) =
delete;
73 auto operator=(AccGenericSycl
const&) -> AccGenericSycl& =
delete;
74 auto operator=(AccGenericSycl&&) -> AccGenericSycl& =
delete;
77 Vec<TDim, TIdx>
const& threadElemExtent,
78 sycl::nd_item<TDim::value> work_item,
79 sycl::local_accessor<std::byte> dyn_shared_acc,
80 sycl::local_accessor<std::byte> st_shared_acc)
81 : WorkDivGenericSycl<TDim, TIdx>{threadElemExtent, work_item}
82 , gb::IdxGbGenericSycl<TDim, TIdx>{work_item}
83 , bt::IdxBtGenericSycl<TDim, TIdx>{work_item}
84 , BlockSharedMemDynGenericSycl{dyn_shared_acc}
85 , BlockSharedMemStGenericSycl{st_shared_acc}
86 , BlockSyncGenericSycl<TDim>{work_item}
87 # ifndef ALPAKA_DISABLE_VENDOR_RNG
88 , rand::RandGenericSycl<TDim>{work_item}
90 , warp::WarpGenericSycl<TDim>{work_item}
99 template<
template<
typename,
typename>
typename TAcc,
typename TDim,
typename TIdx>
100 struct AccType<TAcc<TDim, TIdx>, std::enable_if_t<std::is_base_of_v<AccGenericSycl<TDim, TIdx>, TAcc<TDim, TIdx>>>>
102 using type = TAcc<TDim, TIdx>;
106 template<
template<
typename,
typename>
typename TAcc,
typename TDim,
typename TIdx>
107 struct GetAccDevProps<
109 std::enable_if_t<std::is_base_of_v<AccGenericSycl<TDim, TIdx>, TAcc<TDim, TIdx>>>>
111 static auto getAccDevProps(
typename DevType<TAcc<TDim, TIdx>>::type
const& dev) -> AccDevProps<TDim, TIdx>
113 auto const device = dev.getNativeHandle().first;
114 auto const max_threads_dim
115 = device.template get_info<sycl::info::device::max_work_item_sizes<TDim::value>>();
116 Vec<TDim, TIdx> max_threads_dim_vec{};
117 for(
int i = 0; i < static_cast<int>(TDim::value); i++)
118 max_threads_dim_vec[i] = alpaka::core::clipCast<TIdx>(max_threads_dim[i]);
120 alpaka::core::clipCast<TIdx>(device.template get_info<sycl::info::device::max_compute_units>()),
122 getExtentVecEnd<TDim>(
Vec<DimInt<3u>, TIdx>(
132 alpaka::core::clipCast<TIdx>(device.template get_info<sycl::info::device::max_work_group_size>()),
138 device.template get_info<sycl::info::device::local_mem_size>(),
145 template<
template<
typename,
typename>
typename TAcc,
typename TDim,
typename TIdx>
146 struct DimType<TAcc<TDim, TIdx>, std::enable_if_t<std::is_base_of_v<AccGenericSycl<TDim, TIdx>, TAcc<TDim, TIdx>>>>
152 template<
template<
typename,
typename>
typename TAcc,
typename TDim,
typename TIdx>
153 struct IdxType<TAcc<TDim, TIdx>, std::enable_if_t<std::is_base_of_v<AccGenericSycl<TDim, TIdx>, TAcc<TDim, TIdx>>>>
ALPAKA_NO_HOST_ACC_WARNING static constexpr ALPAKA_FN_HOST_ACC auto all(TVal const &val) -> Vec< TDim, TVal >
Single value constructor.
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto max(T const &max_ctx, Tx const &x, Ty const &y)
Returns the larger of two arguments. NaNs are treated as missing data (between a NaN and a numeric va...
The alpaka accelerator library.
ALPAKA_FN_HOST auto getAccDevProps(TDev const &dev) -> AccDevProps< Dim< TAcc >, Idx< TAcc >>
ALPAKA_FN_HOST auto getMemBytes(TDev const &dev) -> std::size_t
Vec(TFirstIndex &&, TRestIndices &&...) -> Vec< DimInt< 1+sizeof...(TRestIndices)>, std::decay_t< TFirstIndex >>
alpaka::meta::InheritFromList< alpaka::meta::Unique< std::tuple< TGridAtomic, TBlockAtomic, TThreadAtomic, concepts::Implements< ConceptAtomicGrids, TGridAtomic >, concepts::Implements< ConceptAtomicBlocks, TBlockAtomic >, concepts::Implements< ConceptAtomicThreads, TThreadAtomic > >> > AtomicHierarchy
build a single class to inherit from different atomic implementations