27 #ifdef __SYCL_DEVICE_ONLY__
29 #if defined(__NVPTX__)
35 #include <type_traits>
38 inline namespace _V1 {
42 template <
int Dimensions>
45 return id<1>(linear_id);
49 result[0] = linear_id /
r[1];
50 result[1] = linear_id %
r[1];
55 result[0] = linear_id / (
r[1] *
r[2]);
56 result[1] = (linear_id % (
r[1] *
r[2])) /
r[2];
57 result[2] = linear_id %
r[2];
63 auto local_range = g.get_local_range();
64 auto result = local_range[0];
65 for (
size_t i = 1; i < Group::dimensions; ++i)
66 result *= local_range[i];
72 #ifdef __SYCL_DEVICE_ONLY__
73 if constexpr (std::is_same_v<Group,
group<1>> ||
76 auto it = sycl::detail::Builder::getNDItem<Group::dimensions>();
77 return it.get_local_linear_id();
80 return g.get_local_linear_id();
90 template <
typename T,
typename BinaryOperation>
struct is_native_op {
100 template <
typename T,
typename BinaryOperation>
103 std::is_same_v<BinaryOperation, sycl::plus<std::remove_const_t<T>>> ||
104 std::is_same_v<BinaryOperation, sycl::plus<std::add_const_t<T>>> ||
105 std::is_same_v<BinaryOperation, sycl::plus<void>>>;
108 template <
typename T,
typename BinaryOperation>
111 std::is_same_v<BinaryOperation, sycl::multiplies<std::remove_const_t<T>>> ||
112 std::is_same_v<BinaryOperation,
114 std::is_same_v<BinaryOperation, sycl::multiplies<void>>>;
120 template <
typename T,
typename =
void>
124 template <
typename T>
126 std::integral_constant<bool, sycl::detail::is_complex<T>::value ||
127 sycl::detail::is_arithmetic<T>::value>;
129 template <
typename T>
131 : std::bool_constant<is_vec<T>::value &&
132 (is_arithmetic<T>::value ||
133 is_complex<vector_element_t<T>>::value)> {};
136 template <
typename T,
typename BinaryOperation>
140 : std::true_type::value)>;
146 template <
template <
typename>
typename F,
typename T,
int n>
157 template <
typename T>
159 template <
typename T>
167 template <
typename T,
class BinaryOperation>
168 constexpr std::enable_if_t<
174 template <
typename T,
class BinaryOperation>
175 constexpr std::enable_if_t<
176 (is_complex<T>::value && is_multiplies<T, BinaryOperation>::value), T>
181 template <
typename T,
class BinaryOperation>
183 return sycl::known_identity_v<BinaryOperation, T>;
187 template <
typename Group,
typename Ptr,
class Function>
188 Function
for_each(Group g, Ptr first, Ptr last, Function f) {
189 #ifdef __SYCL_DEVICE_ONLY__
192 for (Ptr p = first + offset; p < last; p += stride) {
202 "Group algorithms are not supported on host.");
211 template <
typename Group,
typename T,
class BinaryOperation>
212 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
213 (detail::is_scalar_arithmetic<T>::value ||
214 (detail::is_complex<T>::value &&
215 detail::is_multiplies<T, BinaryOperation>::value)) &&
220 std::is_same_v<decltype(binary_op(
x,
x)), T>,
221 "Result type of binary_op must match reduction accumulation type.");
222 #ifdef __SYCL_DEVICE_ONLY__
223 #if defined(__NVPTX__)
224 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
227 #if (__SYCL_CUDA_ARCH__ >= 800)
228 return detail::masked_reduction_cuda_sm80(g,
x, binary_op, MemberMask[0]);
230 return detail::masked_reduction_cuda_shfls(g,
x, binary_op, MemberMask[0]);
234 return sycl::detail::calc<__spv::GroupOperation::Reduce>(
235 g,
typename sycl::detail::GroupOpTag<T>::type(),
x, binary_op);
239 "Group algorithms are not supported on host.");
245 template <
typename Group,
typename T,
class BinaryOperation>
246 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
247 detail::is_complex<T>::value &&
248 detail::is_native_op<T, sycl::plus<T>>::value &&
249 detail::is_plus<T, BinaryOperation>::value),
252 #ifdef __SYCL_DEVICE_ONLY__
261 "Group algorithms are not supported on host.");
265 template <
typename Group,
typename T,
class BinaryOperation>
266 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
267 detail::is_vector_arithmetic_or_complex<T>::value &&
272 std::is_same_v<decltype(binary_op(
x,
x)), T>,
273 "Result type of binary_op must match reduction accumulation type.");
275 typename detail::get_scalar_binary_op<BinaryOperation>::type
285 template <
typename Group,
typename V,
typename T,
class BinaryOperation>
287 (is_group_v<std::decay_t<Group>> &&
288 (detail::is_scalar_arithmetic<V>::value || detail::is_complex<V>::value) &&
289 (detail::is_scalar_arithmetic<T>::value || detail::is_complex<T>::value) &&
291 detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value &&
292 std::is_convertible_v<V, T>),
296 std::is_same_v<decltype(binary_op(init,
x)), T>,
297 "Result type of binary_op must match reduction accumulation type.");
298 #ifdef __SYCL_DEVICE_ONLY__
303 "Group algorithms are not supported on host.");
307 template <
typename Group,
typename V,
typename T,
class BinaryOperation>
308 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
309 detail::is_vector_arithmetic_or_complex<V>::value &&
310 detail::is_vector_arithmetic_or_complex<T>::value &&
316 std::is_same_v<decltype(binary_op(init,
x)), T>,
317 "Result type of binary_op must match reduction accumulation type.");
318 typename detail::get_scalar_binary_op<BinaryOperation>::type
320 #ifdef __SYCL_DEVICE_ONLY__
322 for (
int s = 0; s <
x.size(); ++s) {
330 "Group algorithms are not supported on host.");
335 template <
typename Group,
typename Ptr,
typename T,
class BinaryOperation>
337 (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr> &&
340 detail::is_arithmetic_or_complex<T>::value &&
341 detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value &&
344 joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
346 std::is_same_v<decltype(binary_op(init, *first)), T>,
347 "Result type of binary_op must match reduction accumulation type.");
348 #ifdef __SYCL_DEVICE_ONLY__
349 T partial = detail::identity_for_ga_op<T, BinaryOperation>();
352 partial = binary_op(partial,
x);
359 "Group algorithms are not supported on host.");
363 template <
typename Group,
typename Ptr,
class BinaryOperation>
365 (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr> &&
371 joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
372 #ifdef __SYCL_DEVICE_ONLY__
374 T init = detail::identity_for_ga_op<T, BinaryOperation>();
382 "Group algorithms are not supported on host.");
387 template <
typename Group>
388 std::enable_if_t<is_group_v<std::decay_t<Group>>,
bool>
390 #ifdef __SYCL_DEVICE_ONLY__
391 #if defined(__NVPTX__)
392 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
397 return sycl::detail::spirv::GroupAny(g, pred);
402 "Group algorithms are not supported on host.");
406 template <
typename Group,
typename T,
class Predicate>
413 template <
typename Group,
typename Ptr,
class Predicate>
414 std::enable_if_t<(is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr>),
417 #ifdef __SYCL_DEVICE_ONLY__
419 bool partial =
false;
428 "Group algorithms are not supported on host.");
433 template <
typename Group>
434 std::enable_if_t<is_group_v<std::decay_t<Group>>,
bool>
436 #ifdef __SYCL_DEVICE_ONLY__
437 #if defined(__NVPTX__)
438 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
443 return sycl::detail::spirv::GroupAll(g, pred);
448 "Group algorithms are not supported on host.");
452 template <
typename Group,
typename T,
class Predicate>
453 std::enable_if_t<is_group_v<std::decay_t<Group>>,
bool>
459 template <
typename Group,
typename Ptr,
class Predicate>
460 std::enable_if_t<(is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr>),
463 #ifdef __SYCL_DEVICE_ONLY__
474 "Group algorithms are not supported on host.");
479 template <
typename Group>
480 std::enable_if_t<is_group_v<std::decay_t<Group>>,
bool>
482 #ifdef __SYCL_DEVICE_ONLY__
483 #if defined(__NVPTX__)
484 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
489 return sycl::detail::spirv::GroupAll(g, !pred);
494 "Group algorithms are not supported on host.");
498 template <
typename Group,
typename T,
class Predicate>
499 std::enable_if_t<is_group_v<std::decay_t<Group>>,
bool>
505 template <
typename Group,
typename Ptr,
class Predicate>
506 std::enable_if_t<(is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr>),
509 #ifdef __SYCL_DEVICE_ONLY__
517 "Group algorithms are not supported on host.");
524 template <
typename Group,
typename T>
525 std::enable_if_t<((std::is_same_v<std::decay_t<Group>, sub_group> ||
527 std::decay_t<Group>>) &&
528 (std::is_trivially_copyable_v<T> ||
529 detail::is_vec<T>::value)),
532 #ifdef __SYCL_DEVICE_ONLY__
533 return sycl::detail::spirv::ShuffleDown(g,
x, delta);
539 "Sub-groups are not supported on host.");
546 template <
typename Group,
typename T>
547 std::enable_if_t<((std::is_same_v<std::decay_t<Group>, sub_group> ||
549 std::decay_t<Group>>) &&
550 (std::is_trivially_copyable_v<T> ||
551 detail::is_vec<T>::value)),
554 #ifdef __SYCL_DEVICE_ONLY__
555 return sycl::detail::spirv::ShuffleUp(g,
x, delta);
561 "Sub-groups are not supported on host.");
568 template <
typename Group,
typename T>
569 std::enable_if_t<((std::is_same_v<std::decay_t<Group>, sub_group> ||
571 std::decay_t<Group>>) &&
572 (std::is_trivially_copyable_v<T> ||
573 detail::is_vec<T>::value)),
576 #ifdef __SYCL_DEVICE_ONLY__
577 return sycl::detail::spirv::ShuffleXor(g,
x, mask);
583 "Sub-groups are not supported on host.");
590 template <
typename Group,
typename T>
591 std::enable_if_t<((std::is_same_v<std::decay_t<Group>, sub_group> ||
593 std::decay_t<Group>>) &&
594 (std::is_trivially_copyable_v<T> ||
595 detail::is_vec<T>::value)),
598 #ifdef __SYCL_DEVICE_ONLY__
599 return sycl::detail::spirv::Shuffle(g,
x, local_id);
605 "Sub-groups are not supported on host.");
612 template <
typename Group,
typename T>
613 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
614 (std::is_trivially_copyable_v<T> ||
615 detail::is_vec<T>::value)),
618 #ifdef __SYCL_DEVICE_ONLY__
619 #if defined(__NVPTX__)
620 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
626 return sycl::detail::spirv::GroupBroadcast(g,
x, local_id);
632 "Group algorithms are not supported on host.");
636 template <
typename Group,
typename T>
637 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
638 (std::is_trivially_copyable_v<T> ||
639 detail::is_vec<T>::value)),
642 #ifdef __SYCL_DEVICE_ONLY__
649 (void)linear_local_id;
651 "Group algorithms are not supported on host.");
655 template <
typename Group,
typename T>
656 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
657 (std::is_trivially_copyable_v<T> ||
658 detail::is_vec<T>::value)),
661 #ifdef __SYCL_DEVICE_ONLY__
667 "Group algorithms are not supported on host.");
675 template <
typename Group,
typename T,
class BinaryOperation>
676 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
677 (detail::is_scalar_arithmetic<T>::value ||
678 (detail::is_complex<T>::value &&
679 detail::is_multiplies<T, BinaryOperation>::value)) &&
683 static_assert(std::is_same_v<decltype(binary_op(
x,
x)), T>,
684 "Result type of binary_op must match scan accumulation type.");
685 #ifdef __SYCL_DEVICE_ONLY__
686 #if defined(__NVPTX__)
687 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
688 return detail::masked_scan_cuda_shfls<__spv::GroupOperation::ExclusiveScan>(
710 auto res = sycl::detail::calc<__spv::GroupOperation::ExclusiveScan>(
711 g,
typename sycl::detail::GroupOpTag<T>::type(),
x, binary_op);
712 if constexpr ((std::is_same_v<signed char, T> ||
713 std::is_same_v<signed short, T> ||
714 (std::is_signed_v<char> && std::is_same_v<char, T>)) &&
715 detail::is_max_or_min<BinaryOperation>::value) {
718 return sycl::known_identity_v<BinaryOperation, T>;
724 "Group algorithms are not supported on host.");
730 template <
typename Group,
typename T,
class BinaryOperation>
731 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
732 detail::is_complex<T>::value &&
733 detail::is_native_op<T, sycl::plus<T>>::value &&
734 detail::is_plus<T, BinaryOperation>::value),
737 #ifdef __SYCL_DEVICE_ONLY__
746 "Group algorithms are not supported on host.");
750 template <
typename Group,
typename T,
class BinaryOperation>
751 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
752 detail::is_vector_arithmetic_or_complex<T>::value &&
756 static_assert(std::is_same_v<decltype(binary_op(
x,
x)), T>,
757 "Result type of binary_op must match scan accumulation type.");
759 typename detail::get_scalar_binary_op<BinaryOperation>::type
761 for (
int s = 0; s <
x.size(); ++s) {
769 template <
typename Group,
typename V,
typename T,
class BinaryOperation>
770 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
771 detail::is_vector_arithmetic_or_complex<V>::value &&
772 detail::is_vector_arithmetic_or_complex<T>::value &&
777 static_assert(std::is_same_v<decltype(binary_op(init,
x)), T>,
778 "Result type of binary_op must match scan accumulation type.");
780 typename detail::get_scalar_binary_op<BinaryOperation>::type
782 for (
int s = 0; s <
x.size(); ++s) {
788 template <
typename Group,
typename V,
typename T,
class BinaryOperation>
790 (is_group_v<std::decay_t<Group>> &&
791 (detail::is_scalar_arithmetic<V>::value || detail::is_complex<V>::value) &&
792 (detail::is_scalar_arithmetic<T>::value || detail::is_complex<T>::value) &&
794 detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value &&
795 std::is_convertible_v<V, T>),
798 static_assert(std::is_same_v<decltype(binary_op(init,
x)), T>,
799 "Result type of binary_op must match scan accumulation type.");
800 #ifdef __SYCL_DEVICE_ONLY__
801 typename Group::linear_id_type local_linear_id =
804 if (local_linear_id == 0) {
805 y = binary_op(init,
y);
808 if (local_linear_id == 0) {
815 "Group algorithms are not supported on host.");
820 template <
typename Group,
typename InPtr,
typename OutPtr,
typename T,
821 class BinaryOperation>
823 (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
824 detail::is_pointer_v<OutPtr> &&
829 detail::is_arithmetic_or_complex<T>::value &&
831 detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value),
834 BinaryOperation binary_op) {
835 static_assert(std::is_same_v<decltype(binary_op(init, *first)), T>,
836 "Result type of binary_op must match scan accumulation type.");
837 #ifdef __SYCL_DEVICE_ONLY__
840 ptrdiff_t N = last - first;
841 auto roundup = [=](
const ptrdiff_t &v,
842 const ptrdiff_t &divisor) -> ptrdiff_t {
843 return ((v + divisor - 1) / divisor) * divisor;
845 typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
848 for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
849 ptrdiff_t i = chunk + offset;
866 "Group algorithms are not supported on host.");
870 template <
typename Group,
typename InPtr,
typename OutPtr,
871 class BinaryOperation>
873 (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
874 detail::is_pointer_v<OutPtr> &&
879 detail::is_native_op<typename detail::remove_pointer<OutPtr>::type,
880 BinaryOperation>::value &&
883 BinaryOperation>::value),
886 BinaryOperation binary_op) {
887 static_assert(std::is_same_v<decltype(binary_op(*first, *first)),
889 "Result type of binary_op must match scan accumulation type.");
891 T init = detail::identity_for_ga_op<T, BinaryOperation>();
899 template <
typename Group,
typename T,
class BinaryOperation>
900 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
901 detail::is_vector_arithmetic_or_complex<T>::value &&
905 static_assert(std::is_same_v<decltype(binary_op(
x,
x)), T>,
906 "Result type of binary_op must match scan accumulation type.");
908 typename detail::get_scalar_binary_op<BinaryOperation>::type
910 for (
int s = 0; s <
x.size(); ++s) {
916 template <
typename Group,
typename T,
class BinaryOperation>
917 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
918 (detail::is_scalar_arithmetic<T>::value ||
919 (detail::is_complex<T>::value &&
920 detail::is_multiplies<T, BinaryOperation>::value)) &&
924 static_assert(std::is_same_v<decltype(binary_op(
x,
x)), T>,
925 "Result type of binary_op must match scan accumulation type.");
926 #ifdef __SYCL_DEVICE_ONLY__
927 #if defined(__NVPTX__)
928 if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
929 return detail::masked_scan_cuda_shfls<__spv::GroupOperation::InclusiveScan>(
934 return sycl::detail::calc<__spv::GroupOperation::InclusiveScan>(
935 g,
typename sycl::detail::GroupOpTag<T>::type(),
x, binary_op);
939 "Group algorithms are not supported on host.");
944 template <
typename Group,
typename T,
class BinaryOperation>
945 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
946 detail::is_complex<T>::value &&
947 detail::is_native_op<T, sycl::plus<T>>::value &&
948 detail::is_plus<T, BinaryOperation>::value),
951 #ifdef __SYCL_DEVICE_ONLY__
960 "Group algorithms are not supported on host.");
966 template <
typename Group,
typename V,
class BinaryOperation,
typename T>
968 (is_group_v<std::decay_t<Group>> &&
969 (detail::is_scalar_arithmetic<V>::value || detail::is_complex<V>::value) &&
970 (detail::is_scalar_arithmetic<T>::value || detail::is_complex<T>::value) &&
972 detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value &&
973 std::is_convertible_v<V, T>),
976 static_assert(std::is_same_v<decltype(binary_op(init,
x)), T>,
977 "Result type of binary_op must match scan accumulation type.");
978 #ifdef __SYCL_DEVICE_ONLY__
981 y = binary_op(init,
y);
987 "Group algorithms are not supported on host.");
991 template <
typename Group,
typename V,
class BinaryOperation,
typename T>
992 std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
993 detail::is_vector_arithmetic_or_complex<V>::value &&
994 detail::is_vector_arithmetic_or_complex<T>::value &&
999 static_assert(std::is_same_v<decltype(binary_op(init,
x)), T>,
1000 "Result type of binary_op must match scan accumulation type.");
1002 typename detail::get_scalar_binary_op<BinaryOperation>::type
1004 for (
int s = 0; s <
x.size(); ++s) {
1011 template <
typename Group,
typename InPtr,
typename OutPtr,
1012 class BinaryOperation,
typename T>
1014 (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
1015 detail::is_pointer_v<OutPtr> &&
1020 detail::is_arithmetic_or_complex<T>::value &&
1022 detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value),
1025 BinaryOperation binary_op, T init) {
1026 static_assert(std::is_same_v<decltype(binary_op(init, *first)), T>,
1027 "Result type of binary_op must match scan accumulation type.");
1028 #ifdef __SYCL_DEVICE_ONLY__
1031 ptrdiff_t N = last - first;
1032 auto roundup = [=](
const ptrdiff_t &v,
1033 const ptrdiff_t &divisor) -> ptrdiff_t {
1034 return ((v + divisor - 1) / divisor) * divisor;
1036 typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
1039 for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
1040 ptrdiff_t i = chunk + offset;
1056 "Group algorithms are not supported on host.");
1060 template <
typename Group,
typename InPtr,
typename OutPtr,
1061 class BinaryOperation>
1063 (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
1064 detail::is_pointer_v<OutPtr> &&
1067 detail::is_native_op<typename detail::remove_pointer<OutPtr>::type,
1068 BinaryOperation>::value &&
1071 BinaryOperation>::value),
1074 BinaryOperation binary_op) {
1075 static_assert(std::is_same_v<decltype(binary_op(*first, *first)),
1077 "Result type of binary_op must match scan accumulation type.");
1080 T init = detail::identity_for_ga_op<T, BinaryOperation>();
A unique identifier of an item in an index space.
class sycl::vec ///////////////////////// Provides a cross-patform vector class template that works e...
id< Dimensions > linear_id_to_id(range< Dimensions >, size_t linear_id)
auto get_local_linear_id(Group g)
std::integral_constant< bool, std::is_same_v< BinaryOperation, sycl::multiplies< std::remove_const_t< T > >>||std::is_same_v< BinaryOperation, sycl::multiplies< std::add_const_t< T > >>||std::is_same_v< BinaryOperation, sycl::multiplies< void > >> is_multiplies
auto get_local_linear_range(Group g)
std::integral_constant< bool, std::is_same_v< BinaryOperation, sycl::plus< std::remove_const_t< T > >>||std::is_same_v< BinaryOperation, sycl::plus< std::add_const_t< T > >>||std::is_same_v< BinaryOperation, sycl::plus< void > >> is_plus
boost::mp11::mp_set_contains< TypeList, std::remove_cv_t< T > > is_contained
uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id)
boost::mp11::mp_list< T... > type_list
std::integral_constant< bool, sycl::detail::is_complex< T >::value||sycl::detail::is_arithmetic< T >::value > is_arithmetic_or_complex
ext::oneapi::sub_group_mask GetMask(NonUniformGroup Group)
std::integral_constant< bool,(is_complex< T >::value ?(is_plus< T, BinaryOperation >::value||is_multiplies< T, BinaryOperation >::value) :std::true_type::value)> is_plus_or_multiplies_if_complex
sycl::vec< unsigned, 4 > ExtractMask(ext::oneapi::sub_group_mask Mask)
id< 3 > linear_id_to_id(range< 3 > r, size_t linear_id)
type_list< sycl::plus< T >, sycl::bit_or< T >, sycl::bit_xor< T >, sycl::bit_and< T >, sycl::maximum< T >, sycl::minimum< T >, sycl::multiplies< T >, sycl::logical_or< T >, sycl::logical_and< T > > native_op_list
constexpr std::enable_if_t<(is_complex< T >::value &&is_plus< T, BinaryOperation >::value), T > identity_for_ga_op()
Function for_each(Group g, Ptr first, Ptr last, Function f)
constexpr bool is_user_constructed_group_v
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&(std::is_trivially_copyable_v< T >||detail::is_vec< T >::value)), T > group_broadcast(Group g, T x, typename Group::id_type local_id)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&(detail::is_scalar_arithmetic< T >::value||(detail::is_complex< T >::value &&detail::is_multiplies< T, BinaryOperation >::value)) &&detail::is_native_op< T, BinaryOperation >::value), T > exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op)
std::bit_and< T > bit_and
std::enable_if_t<((std::is_same_v< std::decay_t< Group >, sub_group >||sycl::ext::oneapi::experimental::is_user_constructed_group_v< std::decay_t< Group >>) &&(std::is_trivially_copyable_v< T >||detail::is_vec< T >::value)), T > shift_group_left(Group g, T x, typename Group::linear_id_type delta=1)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< Ptr > &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< Ptr >::type >::value &&detail::is_arithmetic_or_complex< T >::value &&detail::is_plus_or_multiplies_if_complex< T, BinaryOperation >::value &&detail::is_native_op< T, BinaryOperation >::value), T > joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op)
std::enable_if_t< is_group_v< std::decay_t< Group > >, bool > any_of_group(Group g, bool pred)
std::enable_if_t<((std::is_same_v< std::decay_t< Group >, sub_group >||sycl::ext::oneapi::experimental::is_user_constructed_group_v< std::decay_t< Group >>) &&(std::is_trivially_copyable_v< T >||detail::is_vec< T >::value)), T > shift_group_right(Group g, T x, typename Group::linear_id_type delta=1)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< InPtr > &&detail::is_pointer_v< OutPtr > &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< InPtr >::type >::value &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< OutPtr >::type >::value &&detail::is_arithmetic_or_complex< T >::value &&detail::is_native_op< T, BinaryOperation >::value &&detail::is_plus_or_multiplies_if_complex< T, BinaryOperation >::value), OutPtr > joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op)
std::multiplies< T > multiplies
std::bit_xor< T > bit_xor
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< InPtr > &&detail::is_pointer_v< OutPtr > &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< InPtr >::type >::value &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< OutPtr >::type >::value &&detail::is_arithmetic_or_complex< T >::value &&detail::is_native_op< T, BinaryOperation >::value &&detail::is_plus_or_multiplies_if_complex< T, BinaryOperation >::value), OutPtr > joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&(detail::is_scalar_arithmetic< T >::value||(detail::is_complex< T >::value &&detail::is_multiplies< T, BinaryOperation >::value)) &&detail::is_native_op< T, BinaryOperation >::value), T > reduce_over_group(Group g, T x, BinaryOperation binary_op)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_vector_arithmetic_or_complex< T >::value &&detail::is_native_op< T, BinaryOperation >::value), T > inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< Ptr >), bool > joint_any_of(Group g, Ptr first, Ptr last, Predicate pred)
std::enable_if_t<((std::is_same_v< std::decay_t< Group >, sub_group >||sycl::ext::oneapi::experimental::is_user_constructed_group_v< std::decay_t< Group >>) &&(std::is_trivially_copyable_v< T >||detail::is_vec< T >::value)), T > select_from_group(Group g, T x, typename Group::id_type local_id)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< Ptr >), bool > joint_all_of(Group g, Ptr first, Ptr last, Predicate pred)
std::enable_if_t< is_group_v< std::decay_t< Group > >, bool > none_of_group(Group g, bool pred)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
std::enable_if_t<((std::is_same_v< std::decay_t< Group >, sub_group >||sycl::ext::oneapi::experimental::is_user_constructed_group_v< std::decay_t< Group >>) &&(std::is_trivially_copyable_v< T >||detail::is_vec< T >::value)), T > permute_group_by_xor(Group g, T x, typename Group::linear_id_type mask)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< Ptr >), bool > joint_none_of(Group g, Ptr first, Ptr last, Predicate pred)
std::enable_if_t< is_group_v< std::decay_t< Group > >, bool > all_of_group(Group g, bool pred)
static constexpr bool value
std::remove_cv_t< T > type