23 #include <type_traits>
26 inline namespace _V1 {
30 template <
typename Type,
int NumElements>
class __SYCL_EBO vec;
40 static_assert(is_multi_ptr_v<MultiPtrTy>);
42 using DecoratedPtrTy = decltype(DecoratedPtr);
47 #ifdef __SYCL_DEVICE_ONLY__
50 deduce_AS<DecoratedPtrTy>::value>::type *;
52 using ResultTy = TargetElemTy *;
54 return reinterpret_cast<ResultTy
>(DecoratedPtr);
57 template <
typename T, access::address_space Space>
59 std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
62 template <
typename T, access::address_space Space>
64 std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
67 #ifdef __SYCL_DEVICE_ONLY__
74 return sycl::bit_cast<T>(Ret);
80 using BlockT = SelectBlockT<T>;
81 using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
84 return sycl::bit_cast<typename vec<T, N>::vector_t>(Ret);
90 using BlockT = SelectBlockT<T>;
93 sycl::bit_cast<BlockT>(x));
99 using BlockT = SelectBlockT<T>;
100 using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
103 sycl::bit_cast<VecT>(x));
113 inline multi_ptr<T, Space, IsDecorated>
115 if constexpr (IsDecorated == access::decorated::legacy) {
122 Mptr.get_decorated())};
129 namespace ext::oneapi::this_work_item {
145 #ifdef __SYCL_DEVICE_ONLY__
146 return __spirv_SubgroupLocalInvocationId();
149 "Sub-groups are not supported on host.");
154 #ifdef __SYCL_DEVICE_ONLY__
158 "Sub-groups are not supported on host.");
163 #ifdef __SYCL_DEVICE_ONLY__
164 return __spirv_SubgroupSize();
167 "Sub-groups are not supported on host.");
172 #ifdef __SYCL_DEVICE_ONLY__
173 return __spirv_SubgroupMaxSize();
176 "Sub-groups are not supported on host.");
181 #ifdef __SYCL_DEVICE_ONLY__
182 return __spirv_SubgroupId();
185 "Sub-groups are not supported on host.");
190 #ifdef __SYCL_DEVICE_ONLY__
194 "Sub-groups are not supported on host.");
199 #ifdef __SYCL_DEVICE_ONLY__
200 return __spirv_NumSubgroups();
203 "Sub-groups are not supported on host.");
209 #ifdef __SYCL_DEVICE_ONLY__
211 template <
typename CVT,
typename T = std::remove_cv_t<CVT>>
212 std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value, T>
213 load(CVT *cv_src)
const {
214 T *src =
const_cast<T *
>(cv_src);
216 sycl::detail::deduce_AS<T>::value,
217 sycl::access::decorated::yes>(src));
221 template <
typename CVT,
typename T = std::remove_cv_t<CVT>>
222 std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value, T>
223 load(CVT *cv_src)
const {
224 T *src =
const_cast<T *
>(cv_src);
226 #if defined(__NVPTX__) || defined(__AMDGCN__)
229 auto l = __SYCL_GenericCastToPtrExplicit_ToLocal<T>(src);
233 auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(src);
242 template <
typename CVT,
typename T = std::remove_cv_t<CVT>>
246 "Sub-groups are not supported on host.");
257 #ifdef __SYCL_DEVICE_ONLY__
258 #if defined(__NVPTX__) || defined(__AMDGCN__)
261 return sycl::detail::sub_group::load(src);
266 "Sub-groups are not supported on host.");
277 #ifdef __SYCL_DEVICE_ONLY__
282 "Sub-groups are not supported on host.");
285 #ifdef __SYCL_DEVICE_ONLY__
286 #if defined(__NVPTX__) || defined(__AMDGCN__)
296 for (
int i = 0; i < N; ++i) {
306 N != 1 && N != 3 && N != 16,
308 load(
const multi_ptr<CVT, Space, IsDecorated> cv_src)
const {
309 multi_ptr<T, Space, IsDecorated> src =
311 return sycl::detail::sub_group::load<N, T>(src);
320 load(
const multi_ptr<CVT, Space, IsDecorated> cv_src)
const {
321 multi_ptr<T, Space, IsDecorated> src =
323 return {sycl::detail::sub_group::load<8, T>(src),
324 sycl::detail::sub_group::load<8, T>(src +
334 load(
const multi_ptr<CVT, Space, IsDecorated> cv_src)
const {
335 multi_ptr<T, Space, IsDecorated> src =
338 sycl::detail::sub_group::load<1, T>(src),
348 load(
const multi_ptr<CVT, Space, IsDecorated> cv_src)
const {
349 multi_ptr<T, Space, IsDecorated> src =
351 return sycl::detail::sub_group::load(src);
363 "Sub-groups are not supported on host.");
375 #ifdef __SYCL_DEVICE_ONLY__
377 for (
int i = 0; i < N; ++i) {
384 "Sub-groups are not supported on host.");
388 #ifdef __SYCL_DEVICE_ONLY__
390 template <
typename T>
391 std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value>
394 sycl::detail::deduce_AS<T>::value,
395 sycl::access::decorated::yes>(dst),
400 template <
typename T>
401 std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value>
404 #if defined(__NVPTX__) || defined(__AMDGCN__)
407 auto l = __SYCL_GenericCastToPtrExplicit_ToLocal<T>(dst);
413 auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(dst);
424 template <
typename T>
void store(T *dst,
const T &
x)
const {
428 "Sub-groups are not supported on host.");
437 #ifdef __SYCL_DEVICE_ONLY__
438 #if defined(__NVPTX__) || defined(__AMDGCN__)
441 sycl::detail::sub_group::store(dst,
x);
447 "Sub-groups are not supported on host.");
456 #ifdef __SYCL_DEVICE_ONLY__
462 "Sub-groups are not supported on host.");
466 #ifdef __SYCL_DEVICE_ONLY__
467 #if defined(__NVPTX__) || defined(__AMDGCN__)
473 for (
int i = 0; i < N; ++i) {
482 N != 1 && N != 3 && N != 16>
483 store(multi_ptr<T, Space, DecorateAddress> dst,
const vec<T, N> &
x)
const {
484 sycl::detail::sub_group::store(dst,
x);
492 store(multi_ptr<T, Space, DecorateAddress> dst,
const vec<T, 1> &
x)
const {
493 sycl::detail::sub_group::store(dst,
x);
501 store(multi_ptr<T, Space, DecorateAddress> dst,
const vec<T, 3> &
x)
const {
502 store<1, T, Space, DecorateAddress>(dst,
x.s0());
512 store(multi_ptr<T, Space, DecorateAddress> dst,
const vec<T, 16> &
x)
const {
513 store<8, T, Space, DecorateAddress>(dst,
x.lo());
528 "Sub-groups are not supported on host.");
537 #ifdef __SYCL_DEVICE_ONLY__
538 for (
int i = 0; i < N; ++i) {
545 "Sub-groups are not supported on host.");
551 "Sub-group barrier with no arguments is deprecated."
552 "Use sycl::group_barrier with the sub-group as the argument instead.")
554 #ifdef __SYCL_DEVICE_ONLY__
563 "Sub-groups are not supported on host.");
568 "Sub-group barrier accepting fence_space is deprecated."
569 "Use sycl::group_barrier with the sub-group as the argument instead.")
571 #ifdef __SYCL_DEVICE_ONLY__
578 "Sub-groups are not supported on host.");
583 #ifdef __SYCL_DEVICE_ONLY__
587 "Sub-groups are not supported on host.");
592 #ifdef __SYCL_DEVICE_ONLY__
596 "Sub-groups are not supported on host.");
600 bool leader()
const {
601 #ifdef __SYCL_DEVICE_ONLY__
605 "Sub-groups are not supported on host.");
611 #ifdef __SYCL_DEVICE_ONLY__
612 return lhs.get_group_id() == rhs.get_group_id();
617 "Sub-groups are not supported on host.");
622 #ifdef __SYCL_DEVICE_ONLY__
623 return !(lhs == rhs);
628 "Sub-groups are not supported on host.");
Identifies an instance of the function object executing at each point in an nd_range.
__ESIMD_API void barrier()
Generic work-group barrier.
std::bool_constant<!std::is_same_v< void, SelectBlockT< T > > &&Space==access::address_space::local_space > AcceptableForLocalLoadStore
select_cl_scalar_integral_unsigned_t< T > SelectBlockT
auto convertToBlockPtr(MultiPtrTy MultiPtr)
std::bool_constant<!std::is_same_v< void, SelectBlockT< T > > &&Space==access::address_space::global_space > AcceptableForGlobalLoadStore
auto convertToOpenCLType(T &&x)
auto get_local_linear_range(Group g)
select_apply_cl_scalar_t< T, sycl::opencl::cl_uchar, sycl::opencl::cl_ushort, sycl::opencl::cl_uint, sycl::opencl::cl_ulong > select_cl_scalar_integral_unsigned_t
constexpr __spv::MemorySemanticsMask::Flag getSPIRVMemorySemanticsMask(memory_order)
multi_ptr< T, Space, IsDecorated > GetUnqualMultiPtr(const multi_ptr< CVT, Space, IsDecorated > &Mptr)
bool operator==(const cache_config &lhs, const cache_config &rhs)
sycl::sub_group get_sub_group()
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
typename remove_decoration< T >::type remove_decoration_t
__SYCL_CONVERGENT__ __DPCPP_SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
std::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore< T, Space >::value > store(multi_ptr< T, Space, DecorateAddress > dst, const T &x) const
linear_id_type get_group_linear_id() const
std::enable_if_t< sycl::detail::sub_group::AcceptableForGlobalLoadStore< T, Space >::value > store(multi_ptr< T, Space, DecorateAddress > dst, const vec< T, N > &x) const
__SYCL_DEPRECATED("Sub-group barrier with no arguments is deprecated." "Use sycl::group_barrier with the sub-group as the argument instead.") void barrier() const
range_type get_max_local_range() const
std::enable_if_t< sycl::detail::sub_group::AcceptableForGlobalLoadStore< T, Space >::value, vec< T, N > > load(const multi_ptr< CVT, Space, IsDecorated > src) const
id_type get_group_id() const
range_type get_local_range() const
static constexpr int dimensions
static constexpr sycl::memory_scope fence_scope
__SYCL_DEPRECATED("Sub-group barrier accepting fence_space is deprecated." "Use sycl::group_barrier with the sub-group as the argument instead.") void barrier(access
std::enable_if_t< sycl::detail::sub_group::AcceptableForGlobalLoadStore< T, Space >::value, T > load(const multi_ptr< CVT, Space, IsDecorated > cv_src) const
std::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore< T, Space >::value > store(multi_ptr< T, Space, DecorateAddress > dst, const vec< T, N > &x) const
std::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore< T, Space >::value, T > load(const multi_ptr< CVT, Space, IsDecorated > cv_src) const
id_type get_local_id() const
void store(T *dst, const T &x) const
friend bool operator!=(const sub_group &lhs, const sub_group &rhs)
linear_id_type get_local_linear_id() const
std::enable_if_t< sycl::detail::sub_group::AcceptableForLocalLoadStore< T, Space >::value, vec< T, N > > load(const multi_ptr< CVT, Space, IsDecorated > cv_src) const
range_type get_group_range() const
std::enable_if_t< sycl::detail::sub_group::AcceptableForGlobalLoadStore< T, Space >::value > store(multi_ptr< T, Space, DecorateAddress > dst, const T &x) const