DPC++ Runtime
Runtime libraries for oneAPI DPC++
group_algorithm.hpp
Go to the documentation of this file.
1 //==------------------------ group_algorithm.hpp ---------------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/detail/array.hpp> // for array
12 #include <sycl/detail/helpers.hpp> // for loop
13 #include <sycl/detail/item_base.hpp> // for id, range
14 #include <sycl/detail/type_list.hpp> // for is_contained, type_list
15 #include <sycl/detail/type_traits.hpp> // for remove_pointer, is_pointer
16 #include <sycl/exception.hpp> // for make_error_code, errc, exception
17 #include <sycl/functional.hpp> // for plus, multiplies, maximum
18 #include <sycl/group.hpp> // for group
19 #include <sycl/half_type.hpp> // for half
20 #include <sycl/id.hpp> // for id
21 #include <sycl/known_identity.hpp> // for known_identity_v
22 #include <sycl/nd_item.hpp> // for nd_item
23 #include <sycl/range.hpp> // for range
24 #include <sycl/sub_group.hpp> // for sub_group
25 #include <sycl/types.hpp> // for vec
26 
27 #ifdef __SYCL_DEVICE_ONLY__
29 #if defined(__NVPTX__)
31 #endif
32 #endif
33 
34 #include <stddef.h> // for size_t
35 #include <type_traits> // for enable_if_t, decay_t, integra...
36 
37 namespace sycl {
38 inline namespace _V1 {
39 namespace detail {
40 
41 // ---- linear_id_to_id
42 template <int Dimensions>
44 template <> inline id<1> linear_id_to_id(range<1>, size_t linear_id) {
45  return id<1>(linear_id);
46 }
47 template <> inline id<2> linear_id_to_id(range<2> r, size_t linear_id) {
48  id<2> result;
49  result[0] = linear_id / r[1];
50  result[1] = linear_id % r[1];
51  return result;
52 }
53 template <> inline id<3> linear_id_to_id(range<3> r, size_t linear_id) {
54  id<3> result;
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];
58  return result;
59 }
60 
61 // ---- get_local_linear_range
62 template <typename Group> inline auto get_local_linear_range(Group g) {
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];
67  return result;
68 }
69 
70 // ---- get_local_linear_id
71 template <typename Group> inline auto get_local_linear_id(Group g) {
72 #ifdef __SYCL_DEVICE_ONLY__
73  if constexpr (std::is_same_v<Group, group<1>> ||
74  std::is_same_v<Group, group<2>> ||
75  std::is_same_v<Group, group<3>>) {
76  auto it = sycl::detail::Builder::getNDItem<Group::dimensions>();
77  return it.get_local_linear_id();
78  }
79 #endif // __SYCL_DEVICE_ONLY__
80  return g.get_local_linear_id();
81 }
82 
83 // ---- is_native_op
84 template <typename T>
89 
90 template <typename T, typename BinaryOperation> struct is_native_op {
91  static constexpr bool value =
92  is_contained<BinaryOperation,
94  is_contained<BinaryOperation,
97 };
98 
99 // ---- is_plus
100 template <typename T, typename BinaryOperation>
101 using is_plus = std::integral_constant<
102  bool,
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>>>;
106 
107 // ---- is_multiplies
108 template <typename T, typename BinaryOperation>
109 using is_multiplies = std::integral_constant<
110  bool,
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>>>;
115 
116 // ---- is_complex
117 // Use SFINAE so that the "true" branch could be implemented in
118 // include/sycl/stl_wrappers/complex that would only be available if STL's
119 // <complex> is included by users.
120 template <typename T, typename = void>
121 struct is_complex : public std::false_type {};
122 
123 // ---- is_arithmetic_or_complex
124 template <typename T>
126  std::integral_constant<bool, sycl::detail::is_complex<T>::value ||
127  sycl::detail::is_arithmetic<T>::value>;
128 
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)> {};
134 
135 // ---- is_plus_or_multiplies_if_complex
136 template <typename T, typename BinaryOperation>
137 using is_plus_or_multiplies_if_complex = std::integral_constant<
140  : std::true_type::value)>;
141 
142 // used to transform a vector op to a scalar op;
143 // e.g. sycl::plus<std::vec<T, N>> to sycl::plus<T>
144 template <typename T> struct get_scalar_binary_op;
145 
146 template <template <typename> typename F, typename T, int n>
147 struct get_scalar_binary_op<F<sycl::vec<T, n>>> {
148  using type = F<T>;
149 };
150 
151 template <template <typename> typename F> struct get_scalar_binary_op<F<void>> {
152  using type = F<void>;
153 };
154 
155 // ---- is_max_or_min
156 template <typename T> struct is_max_or_min : std::false_type {};
157 template <typename T>
158 struct is_max_or_min<sycl::maximum<T>> : std::true_type {};
159 template <typename T>
160 struct is_max_or_min<sycl::minimum<T>> : std::true_type {};
161 
162 // ---- identity_for_ga_op
163 // the group algorithms support std::complex, limited to sycl::plus operation
164 // get the correct identity for group algorithm operation.
165 // TODO: identiy_for_ga_op should be replaced with known_identity once the other
166 // callers of known_identity support complex numbers.
167 template <typename T, class BinaryOperation>
168 constexpr std::enable_if_t<
171  return {0, 0};
172 }
173 
174 template <typename T, class BinaryOperation>
175 constexpr std::enable_if_t<
176  (is_complex<T>::value && is_multiplies<T, BinaryOperation>::value), T>
178  return {1, 0};
179 }
180 
181 template <typename T, class BinaryOperation>
182 constexpr std::enable_if_t<!is_complex<T>::value, T> identity_for_ga_op() {
183  return sycl::known_identity_v<BinaryOperation, T>;
184 }
185 
186 // ---- for_each
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__
190  ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
191  ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
192  for (Ptr p = first + offset; p < last; p += stride) {
193  f(*p);
194  }
195  return f;
196 #else
197  (void)g;
198  (void)first;
199  (void)last;
200  (void)f;
202  "Group algorithms are not supported on host.");
203 #endif
204 }
205 } // namespace detail
206 
207 // ---- reduce_over_group
208 // three argument variant is specialized thrice:
209 // scalar arithmetic, complex (plus only), and vector arithmetic
210 
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)) &&
217  T>
218 reduce_over_group(Group g, T x, BinaryOperation binary_op) {
219  static_assert(
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>) {
225  sycl::vec<unsigned, 4> MemberMask =
227 #if (__SYCL_CUDA_ARCH__ >= 800)
228  return detail::masked_reduction_cuda_sm80(g, x, binary_op, MemberMask[0]);
229 #else
230  return detail::masked_reduction_cuda_shfls(g, x, binary_op, MemberMask[0]);
231 #endif
232  }
233 #endif
234  return sycl::detail::calc<__spv::GroupOperation::Reduce>(
235  g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
236 #else
237  (void)g;
239  "Group algorithms are not supported on host.");
240 #endif
241 }
242 
243 // complex specialization. T is std::complex<float> or similar.
244 // binary op is sycl::plus<std::complex<float>>
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),
250  T>
251 reduce_over_group(Group g, T x, BinaryOperation) {
252 #ifdef __SYCL_DEVICE_ONLY__
253  T result;
254  result.real(reduce_over_group(g, x.real(), sycl::plus<>()));
255  result.imag(reduce_over_group(g, x.imag(), sycl::plus<>()));
256  return result;
257 #else
258  (void)g;
259  (void)x;
261  "Group algorithms are not supported on host.");
262 #endif
263 }
264 
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 &&
269  T>
270 reduce_over_group(Group g, T x, BinaryOperation binary_op) {
271  static_assert(
272  std::is_same_v<decltype(binary_op(x, x)), T>,
273  "Result type of binary_op must match reduction accumulation type.");
274  T result;
275  typename detail::get_scalar_binary_op<BinaryOperation>::type
276  scalar_binary_op{};
277  detail::loop<x.size()>([&](size_t s) {
278  result[s] = reduce_over_group(g, x[s], scalar_binary_op);
279  });
280  return result;
281 }
282 
283 // four argument variant of reduce_over_group specialized twice
284 // (scalar arithmetic || complex), and vector_arithmetic
285 template <typename Group, typename V, typename T, class BinaryOperation>
286 std::enable_if_t<
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>),
293  T>
294 reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
295  static_assert(
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__
299  return binary_op(init, reduce_over_group(g, T(x), binary_op));
300 #else
301  (void)g;
303  "Group algorithms are not supported on host.");
304 #endif
305 }
306 
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 &&
313  T>
314 reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
315  static_assert(
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
319  scalar_binary_op{};
320 #ifdef __SYCL_DEVICE_ONLY__
321  T result = init;
322  for (int s = 0; s < x.size(); ++s) {
323  result[s] =
324  scalar_binary_op(init[s], reduce_over_group(g, x[s], scalar_binary_op));
325  }
326  return result;
327 #else
328  (void)g;
330  "Group algorithms are not supported on host.");
331 #endif
332 }
333 
334 // ---- joint_reduce
335 template <typename Group, typename Ptr, typename T, class BinaryOperation>
336 std::enable_if_t<
337  (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr> &&
339  typename detail::remove_pointer<Ptr>::type>::value &&
340  detail::is_arithmetic_or_complex<T>::value &&
341  detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value &&
343  T>
344 joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
345  static_assert(
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>();
351  g, first, last, [&](const typename detail::remove_pointer<Ptr>::type &x) {
352  partial = binary_op(partial, x);
353  });
354  return reduce_over_group(g, partial, init, binary_op);
355 #else
356  (void)g;
357  (void)last;
359  "Group algorithms are not supported on host.");
360 #endif
361 }
362 
363 template <typename Group, typename Ptr, class BinaryOperation>
364 std::enable_if_t<
365  (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<Ptr> &&
367  typename detail::remove_pointer<Ptr>::type>::value &&
369  typename detail::remove_pointer<Ptr>::type, BinaryOperation>::value),
371 joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
372 #ifdef __SYCL_DEVICE_ONLY__
373  using T = typename detail::remove_pointer<Ptr>::type;
374  T init = detail::identity_for_ga_op<T, BinaryOperation>();
375  return joint_reduce(g, first, last, init, binary_op);
376 #else
377  (void)g;
378  (void)first;
379  (void)last;
380  (void)binary_op;
382  "Group algorithms are not supported on host.");
383 #endif
384 }
385 
386 // ---- any_of_group
387 template <typename Group>
388 std::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
389 any_of_group(Group g, bool pred) {
390 #ifdef __SYCL_DEVICE_ONLY__
391 #if defined(__NVPTX__)
392  if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
393  return __nvvm_vote_any_sync(detail::ExtractMask(detail::GetMask(g))[0],
394  pred);
395  }
396 #endif
397  return sycl::detail::spirv::GroupAny(g, pred);
398 #else
399  (void)g;
400  (void)pred;
402  "Group algorithms are not supported on host.");
403 #endif
404 }
405 
406 template <typename Group, typename T, class Predicate>
407 std::enable_if_t<is_group_v<Group>, bool> any_of_group(Group g, T x,
408  Predicate pred) {
409  return any_of_group(g, pred(x));
410 }
411 
412 // ---- joint_any_of
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>),
415  bool>
416 joint_any_of(Group g, Ptr first, Ptr last, Predicate pred) {
417 #ifdef __SYCL_DEVICE_ONLY__
418  using T = typename detail::remove_pointer<Ptr>::type;
419  bool partial = false;
420  sycl::detail::for_each(g, first, last, [&](T &x) { partial |= pred(x); });
421  return any_of_group(g, partial);
422 #else
423  (void)g;
424  (void)first;
425  (void)last;
426  (void)pred;
428  "Group algorithms are not supported on host.");
429 #endif
430 }
431 
432 // ---- all_of_group
433 template <typename Group>
434 std::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
435 all_of_group(Group g, bool pred) {
436 #ifdef __SYCL_DEVICE_ONLY__
437 #if defined(__NVPTX__)
438  if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
439  return __nvvm_vote_all_sync(detail::ExtractMask(detail::GetMask(g))[0],
440  pred);
441  }
442 #endif
443  return sycl::detail::spirv::GroupAll(g, pred);
444 #else
445  (void)g;
446  (void)pred;
448  "Group algorithms are not supported on host.");
449 #endif
450 }
451 
452 template <typename Group, typename T, class Predicate>
453 std::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
454 all_of_group(Group g, T x, Predicate pred) {
455  return all_of_group(g, pred(x));
456 }
457 
458 // ---- joint_all_of
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>),
461  bool>
462 joint_all_of(Group g, Ptr first, Ptr last, Predicate pred) {
463 #ifdef __SYCL_DEVICE_ONLY__
464  using T = typename detail::remove_pointer<Ptr>::type;
465  bool partial = true;
466  sycl::detail::for_each(g, first, last, [&](T &x) { partial &= pred(x); });
467  return all_of_group(g, partial);
468 #else
469  (void)g;
470  (void)first;
471  (void)last;
472  (void)pred;
474  "Group algorithms are not supported on host.");
475 #endif
476 }
477 
478 // ---- none_of_group
479 template <typename Group>
480 std::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
481 none_of_group(Group g, bool pred) {
482 #ifdef __SYCL_DEVICE_ONLY__
483 #if defined(__NVPTX__)
484  if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
485  return __nvvm_vote_all_sync(detail::ExtractMask(detail::GetMask(g))[0],
486  !pred);
487  }
488 #endif
489  return sycl::detail::spirv::GroupAll(g, !pred);
490 #else
491  (void)g;
492  (void)pred;
494  "Group algorithms are not supported on host.");
495 #endif
496 }
497 
498 template <typename Group, typename T, class Predicate>
499 std::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
500 none_of_group(Group g, T x, Predicate pred) {
501  return none_of_group(g, pred(x));
502 }
503 
504 // ---- joint_none_of
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>),
507  bool>
508 joint_none_of(Group g, Ptr first, Ptr last, Predicate pred) {
509 #ifdef __SYCL_DEVICE_ONLY__
510  return !joint_any_of(g, first, last, pred);
511 #else
512  (void)g;
513  (void)first;
514  (void)last;
515  (void)pred;
517  "Group algorithms are not supported on host.");
518 #endif
519 }
520 
521 // ---- shift_group_left
522 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
523 // copyable.
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)),
530  T>
531 shift_group_left(Group g, T x, typename Group::linear_id_type delta = 1) {
532 #ifdef __SYCL_DEVICE_ONLY__
533  return sycl::detail::spirv::ShuffleDown(g, x, delta);
534 #else
535  (void)g;
536  (void)x;
537  (void)delta;
539  "Sub-groups are not supported on host.");
540 #endif
541 }
542 
543 // ---- shift_group_right
544 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
545 // copyable.
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)),
552  T>
553 shift_group_right(Group g, T x, typename Group::linear_id_type delta = 1) {
554 #ifdef __SYCL_DEVICE_ONLY__
555  return sycl::detail::spirv::ShuffleUp(g, x, delta);
556 #else
557  (void)g;
558  (void)x;
559  (void)delta;
561  "Sub-groups are not supported on host.");
562 #endif
563 }
564 
565 // ---- permute_group_by_xor
566 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
567 // copyable.
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)),
574  T>
575 permute_group_by_xor(Group g, T x, typename Group::linear_id_type mask) {
576 #ifdef __SYCL_DEVICE_ONLY__
577  return sycl::detail::spirv::ShuffleXor(g, x, mask);
578 #else
579  (void)g;
580  (void)x;
581  (void)mask;
583  "Sub-groups are not supported on host.");
584 #endif
585 }
586 
587 // ---- select_from_group
588 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
589 // copyable.
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)),
596  T>
597 select_from_group(Group g, T x, typename Group::id_type local_id) {
598 #ifdef __SYCL_DEVICE_ONLY__
599  return sycl::detail::spirv::Shuffle(g, x, local_id);
600 #else
601  (void)g;
602  (void)x;
603  (void)local_id;
605  "Sub-groups are not supported on host.");
606 #endif
607 }
608 
609 // ---- group_broadcast
610 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
611 // copyable.
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)),
616  T>
617 group_broadcast(Group g, T x, typename Group::id_type local_id) {
618 #ifdef __SYCL_DEVICE_ONLY__
619 #if defined(__NVPTX__)
620  if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
621  auto LocalId = detail::IdToMaskPosition(g, local_id);
622  return __nvvm_shfl_sync_idx_i32(detail::ExtractMask(detail::GetMask(g))[0],
623  x, LocalId, 31);
624  }
625 #endif
626  return sycl::detail::spirv::GroupBroadcast(g, x, local_id);
627 #else
628  (void)g;
629  (void)x;
630  (void)local_id;
632  "Group algorithms are not supported on host.");
633 #endif
634 }
635 
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)),
640  T>
641 group_broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
642 #ifdef __SYCL_DEVICE_ONLY__
643  return group_broadcast(
644  g, x,
645  sycl::detail::linear_id_to_id(g.get_local_range(), linear_local_id));
646 #else
647  (void)g;
648  (void)x;
649  (void)linear_local_id;
651  "Group algorithms are not supported on host.");
652 #endif
653 }
654 
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)),
659  T>
660 group_broadcast(Group g, T x) {
661 #ifdef __SYCL_DEVICE_ONLY__
662  return group_broadcast(g, x, 0);
663 #else
664  (void)g;
665  (void)x;
667  "Group algorithms are not supported on host.");
668 #endif
669 }
670 
671 // ---- exclusive_scan_over_group
672 // this function has two overloads, one with three arguments and one with four
673 // the three argument version is specialized thrice: scalar, complex, and
674 // vector
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)) &&
681  T>
682 exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
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>(
689  g, x, binary_op,
691  }
692 #endif
693  // For the first work item in the group, we cannot return the result
694  // of calc when T is a signed char or short type and the
695  // BinaryOperation is maximum or minimum. calc uses SPIRV group
696  // collective instructions, which only operate on 32 or 64 bit
697  // integers. So, when using calc with a short or char type, the
698  // argument is converted to a 32 bit integer, the 32 bit group
699  // operation is performed, and then converted back to the original
700  // short or char type. For an exclusive scan, the first work item
701  // returns the identity for the supplied operation. However, the
702  // identity of a 32 bit signed integer maximum or minimum when
703  // converted to a signed char or short does not correspond to the
704  // identity of a signed char or short maximum or minimum. For
705  // example, the identity of a signed 32 bit maximum is
706  // INT_MIN=-2**31, and when converted to a signed char, results in
707  // 0. However, the identity of a signed char maximum is
708  // SCHAR_MIN=-2**7. Therefore, we need the following check to
709  // circumvent this issue.
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) {
716  auto local_id = sycl::detail::get_local_linear_id(g);
717  if (local_id == 0)
718  return sycl::known_identity_v<BinaryOperation, T>;
719  }
720  return res;
721 #else
722  (void)g;
724  "Group algorithms are not supported on host.");
725 #endif
726 }
727 
728 // complex specialization. T is std::complex<float> or similar.
729 // binary op is sycl::plus<std::complex<float>>
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),
735  T>
736 exclusive_scan_over_group(Group g, T x, BinaryOperation) {
737 #ifdef __SYCL_DEVICE_ONLY__
738  T result;
739  result.real(exclusive_scan_over_group(g, x.real(), sycl::plus<>()));
740  result.imag(exclusive_scan_over_group(g, x.imag(), sycl::plus<>()));
741  return result;
742 #else
743  (void)g;
744  (void)x;
746  "Group algorithms are not supported on host.");
747 #endif
748 }
749 
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 &&
754  T>
755 exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
756  static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
757  "Result type of binary_op must match scan accumulation type.");
758  T result;
759  typename detail::get_scalar_binary_op<BinaryOperation>::type
760  scalar_binary_op{};
761  for (int s = 0; s < x.size(); ++s) {
762  result[s] = exclusive_scan_over_group(g, x[s], scalar_binary_op);
763  }
764  return result;
765 }
766 
767 // four argument version of exclusive_scan_over_group is specialized twice
768 // once for vector_arithmetic, once for (scalar_arithmetic || complex)
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 &&
775  T>
776 exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
777  static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
778  "Result type of binary_op must match scan accumulation type.");
779  T result;
780  typename detail::get_scalar_binary_op<BinaryOperation>::type
781  scalar_binary_op{};
782  for (int s = 0; s < x.size(); ++s) {
783  result[s] = exclusive_scan_over_group(g, x[s], init[s], scalar_binary_op);
784  }
785  return result;
786 }
787 
788 template <typename Group, typename V, typename T, class BinaryOperation>
789 std::enable_if_t<
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>),
796  T>
797 exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
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 =
803  T y = x;
804  if (local_linear_id == 0) {
805  y = binary_op(init, y);
806  }
807  T scan = exclusive_scan_over_group(g, y, binary_op);
808  if (local_linear_id == 0) {
809  scan = init;
810  }
811  return scan;
812 #else
813  (void)g;
815  "Group algorithms are not supported on host.");
816 #endif
817 }
818 
819 // ---- joint_exclusive_scan
820 template <typename Group, typename InPtr, typename OutPtr, typename T,
821  class BinaryOperation>
822 std::enable_if_t<
823  (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
824  detail::is_pointer_v<OutPtr> &&
826  typename detail::remove_pointer<InPtr>::type>::value &&
828  typename detail::remove_pointer<OutPtr>::type>::value &&
829  detail::is_arithmetic_or_complex<T>::value &&
831  detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value),
832  OutPtr>
833 joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
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__
838  ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
839  ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
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;
844  };
845  typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
846  x = {};
847  T carry = init;
848  for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
849  ptrdiff_t i = chunk + offset;
850  if (i < N) {
851  x = first[i];
852  }
853  T out = exclusive_scan_over_group(g, x, carry, binary_op);
854  if (i < N) {
855  result[i] = out;
856  }
857  carry = group_broadcast(g, binary_op(out, x), stride - 1);
858  }
859  return result + N;
860 #else
861  (void)g;
862  (void)last;
863  (void)result;
864  (void)init;
866  "Group algorithms are not supported on host.");
867 #endif
868 }
869 
870 template <typename Group, typename InPtr, typename OutPtr,
871  class BinaryOperation>
872 std::enable_if_t<
873  (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
874  detail::is_pointer_v<OutPtr> &&
876  typename detail::remove_pointer<InPtr>::type>::value &&
878  typename detail::remove_pointer<OutPtr>::type>::value &&
879  detail::is_native_op<typename detail::remove_pointer<OutPtr>::type,
880  BinaryOperation>::value &&
883  BinaryOperation>::value),
884  OutPtr>
885 joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
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.");
890  using T = typename detail::remove_pointer<OutPtr>::type;
891  T init = detail::identity_for_ga_op<T, BinaryOperation>();
892  return joint_exclusive_scan(g, first, last, result, init, binary_op);
893 }
894 
895 // ---- inclusive_scan_over_group
896 // this function has two overloads, one with three arguments and one with four
897 // the three argument version is specialized thrice: vector, scalar, and
898 // complex
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 &&
903  T>
904 inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
905  static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
906  "Result type of binary_op must match scan accumulation type.");
907  T result;
908  typename detail::get_scalar_binary_op<BinaryOperation>::type
909  scalar_binary_op{};
910  for (int s = 0; s < x.size(); ++s) {
911  result[s] = inclusive_scan_over_group(g, x[s], scalar_binary_op);
912  }
913  return result;
914 }
915 
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)) &&
922  T>
923 inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
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>(
930  g, x, binary_op,
932  }
933 #endif
934  return sycl::detail::calc<__spv::GroupOperation::InclusiveScan>(
935  g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
936 #else
937  (void)g;
939  "Group algorithms are not supported on host.");
940 #endif
941 }
942 
943 // complex specializaiton
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),
949  T>
950 inclusive_scan_over_group(Group g, T x, BinaryOperation) {
951 #ifdef __SYCL_DEVICE_ONLY__
952  T result;
953  result.real(inclusive_scan_over_group(g, x.real(), sycl::plus<>()));
954  result.imag(inclusive_scan_over_group(g, x.imag(), sycl::plus<>()));
955  return result;
956 #else
957  (void)g;
958  (void)x;
960  "Group algorithms are not supported on host.");
961 #endif
962 }
963 
964 // four argument version of inclusive_scan_over_group is specialized twice
965 // once for (scalar_arithmetic || complex) and once for vector_arithmetic
966 template <typename Group, typename V, class BinaryOperation, typename T>
967 std::enable_if_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>),
974  T>
975 inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
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__
979  T y = x;
980  if (sycl::detail::get_local_linear_id(g) == 0) {
981  y = binary_op(init, y);
982  }
983  return inclusive_scan_over_group(g, y, binary_op);
984 #else
985  (void)g;
987  "Group algorithms are not supported on host.");
988 #endif
989 }
990 
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 &&
997  T>
998 inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
999  static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
1000  "Result type of binary_op must match scan accumulation type.");
1001  T result;
1002  typename detail::get_scalar_binary_op<BinaryOperation>::type
1003  scalar_binary_op{};
1004  for (int s = 0; s < x.size(); ++s) {
1005  result[s] = inclusive_scan_over_group(g, x[s], scalar_binary_op, init[s]);
1006  }
1007  return result;
1008 }
1009 
1010 // ---- joint_inclusive_scan
1011 template <typename Group, typename InPtr, typename OutPtr,
1012  class BinaryOperation, typename T>
1013 std::enable_if_t<
1014  (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
1015  detail::is_pointer_v<OutPtr> &&
1017  typename detail::remove_pointer<InPtr>::type>::value &&
1019  typename detail::remove_pointer<OutPtr>::type>::value &&
1020  detail::is_arithmetic_or_complex<T>::value &&
1022  detail::is_plus_or_multiplies_if_complex<T, BinaryOperation>::value),
1023  OutPtr>
1024 joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
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__
1029  ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
1030  ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
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;
1035  };
1036  typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
1037  x = {};
1038  T carry = init;
1039  for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
1040  ptrdiff_t i = chunk + offset;
1041  if (i < N) {
1042  x = first[i];
1043  }
1044  T out = inclusive_scan_over_group(g, x, binary_op, carry);
1045  if (i < N) {
1046  result[i] = out;
1047  }
1048  carry = group_broadcast(g, out, stride - 1);
1049  }
1050  return result + N;
1051 #else
1052  (void)g;
1053  (void)last;
1054  (void)result;
1056  "Group algorithms are not supported on host.");
1057 #endif
1058 }
1059 
1060 template <typename Group, typename InPtr, typename OutPtr,
1061  class BinaryOperation>
1062 std::enable_if_t<
1063  (is_group_v<std::decay_t<Group>> && detail::is_pointer_v<InPtr> &&
1064  detail::is_pointer_v<OutPtr> &&
1066  typename detail::remove_pointer<InPtr>::type>::value &&
1067  detail::is_native_op<typename detail::remove_pointer<OutPtr>::type,
1068  BinaryOperation>::value &&
1071  BinaryOperation>::value),
1072  OutPtr>
1073 joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
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.");
1078 
1079  using T = typename detail::remove_pointer<OutPtr>::type;
1080  T init = detail::identity_for_ga_op<T, BinaryOperation>();
1081  return joint_inclusive_scan(g, first, last, result, binary_op, init);
1082 }
1083 
1084 } // namespace _V1
1085 } // namespace sycl
A unique identifier of an item in an index space.
Definition: id.hpp:36
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: types.hpp:284
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
Definition: type_list.hpp:32
uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id)
boost::mp11::mp_list< T... > type_list
Definition: type_list.hpp:22
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)
void loop(F &&f)
Definition: helpers.hpp:253
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)
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
Definition: functional.hpp:20
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
Definition: functional.hpp:19
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
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::plus< T > plus
Definition: functional.hpp:18
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::bit_or< T > bit_or
Definition: functional.hpp:21
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()
Definition: exception.cpp:87
autodecltype(x) x
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)
Definition: access.hpp:18