DPC++ Runtime
Runtime libraries for oneAPI DPC++
group_algorithm.hpp
Go to the documentation of this file.
1 //==----------- group_algorithm.hpp --- SYCL group algorithm----------------==//
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 #include <CL/__spirv/spirv_ops.hpp>
13 #include <sycl/detail/spirv.hpp>
18 #include <sycl/group.hpp>
19 #include <sycl/group_algorithm.hpp>
20 #include <sycl/nd_item.hpp>
21 
22 namespace sycl {
24 namespace ext::oneapi {
25 
26 // EnableIf shorthands for algorithms that depend only on type
27 template <typename T>
29  sycl::detail::enable_if_t<sycl::detail::is_scalar_arithmetic<T>::value, T>;
30 
31 template <typename T>
33  sycl::detail::enable_if_t<sycl::detail::is_vector_arithmetic<T>::value, T>;
34 
35 template <typename Ptr, typename T>
37  sycl::detail::enable_if_t<sycl::detail::is_pointer<Ptr>::value, T>;
38 
39 template <typename T>
41  sycl::detail::enable_if_t<std::is_trivially_copyable<T>::value &&
42  !sycl::detail::is_vector_arithmetic<T>::value,
43  T>;
44 
45 // EnableIf shorthands for algorithms that depend on type and an operator
46 template <typename T, typename BinaryOperation>
48  sycl::detail::is_scalar_arithmetic<T>::value &&
49  sycl::detail::is_native_op<T, BinaryOperation>::value,
50  T>;
51 
52 template <typename T, typename BinaryOperation>
54  sycl::detail::is_vector_arithmetic<T>::value &&
55  sycl::detail::is_native_op<T, BinaryOperation>::value,
56  T>;
57 
58 // TODO: Lift TriviallyCopyable restriction eventually
59 template <typename T, typename BinaryOperation>
61  (!sycl::detail::is_scalar_arithmetic<T>::value &&
62  !sycl::detail::is_vector_arithmetic<T>::value &&
63  std::is_trivially_copyable<T>::value) ||
64  !sycl::detail::is_native_op<T, BinaryOperation>::value,
65  T>;
66 
67 template <typename Group>
69  "ext::oneapi::all_of is deprecated. Use all_of_group instead.")
70 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> all_of(
71  Group g, bool pred) {
72  return all_of_group(g, pred);
73 }
74 
75 template <typename Group, typename T, class Predicate>
77  "ext::oneapi::all_of is deprecated. Use all_of_group instead.")
78 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> all_of(
79  Group g, T x, Predicate pred) {
80  return all_of_group(g, pred(x));
81 }
82 
83 template <typename Group, typename Ptr, class Predicate>
85  "ext::oneapi::all_of is deprecated. Use joint_all_of instead.")
86 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
87  detail::is_pointer<Ptr>::value),
88  bool> all_of(Group g, Ptr first, Ptr last, Predicate pred) {
89  return joint_all_of(g, first, last, pred);
90 }
91 
92 template <typename Group>
94  "ext::oneapi::any_of is deprecated. Use any_of_group instead.")
95 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> any_of(
96  Group g, bool pred) {
97  return any_of_group(g, pred);
98 }
99 
100 template <typename Group, typename T, class Predicate>
102  "ext::oneapi::any_of is deprecated. Use any_of_group instead.")
103 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> any_of(
104  Group g, T x, Predicate pred) {
105  return any_of_group(g, pred(x));
106 }
107 
108 template <typename Group, typename Ptr, class Predicate>
110  "ext::oneapi::any_of is deprecated. Use joint_any_of instead.")
111 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
112  detail::is_pointer<Ptr>::value),
113  bool> any_of(Group g, Ptr first, Ptr last, Predicate pred) {
114  return joint_any_of(g, first, last, pred);
115 }
116 
117 template <typename Group>
119  "ext::oneapi::none_of is deprecated. Use none_of_group instead.")
120 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> none_of(
121  Group g, bool pred) {
122  return none_of_group(g, pred);
123 }
124 
125 template <typename Group, typename T, class Predicate>
127  "ext::oneapi::none_of is deprecated. Use none_of_group instead.")
128 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> none_of(
129  Group g, T x, Predicate pred) {
130  return none_of_group(g, pred(x));
131 }
132 
133 template <typename Group, typename Ptr, class Predicate>
135  "ext::oneapi::none_of is deprecated. Use joint_none_of instead.")
136 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
137  detail::is_pointer<Ptr>::value),
138  bool> none_of(Group g, Ptr first, Ptr last,
139  Predicate pred) {
140  return joint_none_of(g, first, last, pred);
141 }
142 
143 template <typename Group, typename T>
145  "ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
146 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
147  std::is_trivially_copyable<T>::value &&
148  !detail::is_vector_arithmetic<T>::value),
149  T> broadcast(Group, T x, typename Group::id_type local_id) {
150 #ifdef __SYCL_DEVICE_ONLY__
151  return sycl::detail::spirv::GroupBroadcast<Group>(x, local_id);
152 #else
153  (void)x;
154  (void)local_id;
155  throw runtime_error("Group algorithms are not supported on host device.",
156  PI_ERROR_INVALID_DEVICE);
157 #endif
158 }
159 
160 template <typename Group, typename T>
162  "ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
163 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
164  detail::is_vector_arithmetic<T>::value),
165  T> broadcast(Group g, T x,
166  typename Group::id_type local_id) {
167 #ifdef __SYCL_DEVICE_ONLY__
168  T result;
169  for (int s = 0; s < x.size(); ++s) {
170  result[s] = broadcast(g, x[s], local_id);
171  }
172  return result;
173 #else
174  (void)g;
175  (void)x;
176  (void)local_id;
177  throw runtime_error("Group algorithms are not supported on host device.",
178  PI_ERROR_INVALID_DEVICE);
179 #endif
180 }
181 
182 template <typename Group, typename T>
184  "ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
185 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
186  std::is_trivially_copyable<T>::value &&
187  !detail::is_vector_arithmetic<T>::value),
188  T> broadcast(Group g, T x,
189  typename Group::linear_id_type
190  linear_local_id) {
191 #ifdef __SYCL_DEVICE_ONLY__
192  return broadcast(
193  g, x,
194  sycl::detail::linear_id_to_id(g.get_local_range(), linear_local_id));
195 #else
196  (void)g;
197  (void)x;
198  (void)linear_local_id;
199  throw runtime_error("Group algorithms are not supported on host device.",
200  PI_ERROR_INVALID_DEVICE);
201 #endif
202 }
203 
204 template <typename Group, typename T>
206  "ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
207 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
208  detail::is_vector_arithmetic<T>::value),
209  T> broadcast(Group g, T x,
210  typename Group::linear_id_type
211  linear_local_id) {
212 #ifdef __SYCL_DEVICE_ONLY__
213  T result;
214  for (int s = 0; s < x.size(); ++s) {
215  result[s] = broadcast(g, x[s], linear_local_id);
216  }
217  return result;
218 #else
219  (void)g;
220  (void)x;
221  (void)linear_local_id;
222  throw runtime_error("Group algorithms are not supported on host device.",
223  PI_ERROR_INVALID_DEVICE);
224 #endif
225 }
226 
227 template <typename Group, typename T>
229  "ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
230 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
231  std::is_trivially_copyable<T>::value &&
232  !detail::is_vector_arithmetic<T>::value),
233  T> broadcast(Group g, T x) {
234 #ifdef __SYCL_DEVICE_ONLY__
235  return broadcast(g, x, 0);
236 #else
237  (void)g;
238  (void)x;
239  throw runtime_error("Group algorithms are not supported on host device.",
240  PI_ERROR_INVALID_DEVICE);
241 #endif
242 }
243 
244 template <typename Group, typename T>
246  "ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
247 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
248  detail::is_vector_arithmetic<T>::value),
249  T> broadcast(Group g, T x) {
250 #ifdef __SYCL_DEVICE_ONLY__
251  T result;
252  for (int s = 0; s < x.size(); ++s) {
253  result[s] = broadcast(g, x[s]);
254  }
255  return result;
256 #else
257  (void)g;
258  (void)x;
259  throw runtime_error("Group algorithms are not supported on host device.",
260  PI_ERROR_INVALID_DEVICE);
261 #endif
262 }
263 
264 template <typename Group, typename T, class BinaryOperation>
266  "ext::oneapi::reduce is deprecated. Use reduce_over_group instead.")
267 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
268  detail::is_scalar_arithmetic<T>::value &&
269  detail::is_native_op<T, BinaryOperation>::value),
270  T> reduce(Group g, T x, BinaryOperation binary_op) {
271  return reduce_over_group(g, x, binary_op);
272 }
273 
274 template <typename Group, typename T, class BinaryOperation>
276  "ext::oneapi::reduce is deprecated. Use reduce_over_group instead.")
277 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
278  detail::is_vector_arithmetic<T>::value &&
279  detail::is_native_op<T, BinaryOperation>::value),
280  T> reduce(Group g, T x, BinaryOperation binary_op) {
281  return reduce_over_group(g, x, binary_op);
282 }
283 
284 template <typename Group, typename T, class BinaryOperation>
286  "ext::oneapi::reduce is deprecated. Use reduce_over_group instead.")
287 detail::enable_if_t<(detail::is_sub_group<Group>::value &&
288  std::is_trivially_copyable<T>::value &&
289  (!detail::is_arithmetic<T>::value ||
290  !detail::is_native_op<T, BinaryOperation>::value)),
291  T> reduce(Group g, T x, BinaryOperation op) {
292  T result = x;
293  for (int mask = 1; mask < g.get_max_local_range()[0]; mask *= 2) {
294  T tmp = g.shuffle_xor(result, id<1>(mask));
295  if ((g.get_local_id()[0] ^ mask) < g.get_local_range()[0]) {
296  result = op(result, tmp);
297  }
298  }
299  return g.shuffle(result, 0);
300 }
301 
302 template <typename Group, typename V, typename T, class BinaryOperation>
304  "ext::oneapi::reduce is deprecated. Use reduce_over_group instead.")
305 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
306  detail::is_scalar_arithmetic<V>::value &&
307  detail::is_scalar_arithmetic<T>::value &&
308  detail::is_native_op<V, BinaryOperation>::value &&
309  detail::is_native_op<T, BinaryOperation>::value),
310  T> reduce(Group g, V x, T init, BinaryOperation binary_op) {
311  return reduce_over_group(g, x, init, binary_op);
312 }
313 
314 template <typename Group, typename V, typename T, class BinaryOperation>
316  "ext::oneapi::reduce is deprecated. Use reduce_over_group instead.")
317 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
318  detail::is_vector_arithmetic<V>::value &&
319  detail::is_vector_arithmetic<T>::value &&
320  detail::is_native_op<V, BinaryOperation>::value &&
321  detail::is_native_op<T, BinaryOperation>::value),
322  T> reduce(Group g, V x, T init, BinaryOperation binary_op) {
323  return reduce_over_group(g, x, init, binary_op);
324 }
325 
326 template <typename Group, typename V, typename T, class BinaryOperation>
328  "ext::oneapi::reduce is deprecated. Use reduce_over_group instead.")
329 detail::enable_if_t<(detail::is_sub_group<Group>::value &&
330  std::is_trivially_copyable<T>::value &&
331  std::is_trivially_copyable<V>::value &&
332  (!detail::is_arithmetic<T>::value ||
333  !detail::is_arithmetic<V>::value ||
334  !detail::is_native_op<T, BinaryOperation>::value)),
335  T> reduce(Group g, V x, T init, BinaryOperation op) {
336  T result = x;
337  for (int mask = 1; mask < g.get_max_local_range()[0]; mask *= 2) {
338  T tmp = g.shuffle_xor(result, id<1>(mask));
339  if ((g.get_local_id()[0] ^ mask) < g.get_local_range()[0]) {
340  result = op(result, tmp);
341  }
342  }
343  return g.shuffle(op(init, result), 0);
344 }
345 
346 template <typename Group, typename Ptr, class BinaryOperation>
348  "ext::oneapi::reduce is deprecated. Use joint_reduce instead.")
349 detail::enable_if_t<
350  (detail::is_generic_group<Group>::value && detail::is_pointer<Ptr>::value &&
351  detail::is_arithmetic<typename detail::remove_pointer<Ptr>::type>::value),
352  typename detail::remove_pointer<Ptr>::type> reduce(Group g, Ptr first,
353  Ptr last,
354  BinaryOperation
355  binary_op) {
356  return joint_reduce(g, first, last, binary_op);
357 }
358 
359 template <typename Group, typename Ptr, typename T, class BinaryOperation>
361  "ext::oneapi::reduce is deprecated. Use joint_reduce instead.")
362 detail::enable_if_t<
363  (detail::is_generic_group<Group>::value && detail::is_pointer<Ptr>::value &&
364  detail::is_arithmetic<typename detail::remove_pointer<Ptr>::type>::value &&
365  detail::is_arithmetic<T>::value &&
366  detail::is_native_op<typename detail::remove_pointer<Ptr>::type,
367  BinaryOperation>::value &&
368  detail::is_native_op<T, BinaryOperation>::value),
369  T> reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
370  return joint_reduce(g, first, last, init, binary_op);
371 }
372 
373 template <typename Group, typename T, class BinaryOperation>
374 __SYCL2020_DEPRECATED("ext::oneapi::exclusive_scan is deprecated. Use "
375  "exclusive_scan_over_group instead.")
376 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
377  detail::is_scalar_arithmetic<T>::value &&
378  detail::is_native_op<T, BinaryOperation>::value),
379  T> exclusive_scan(Group g, T x, BinaryOperation binary_op) {
380  return exclusive_scan_over_group(g, x, binary_op);
381 }
382 
383 template <typename Group, typename T, class BinaryOperation>
384 __SYCL2020_DEPRECATED("ext::oneapi::exclusive_scan is deprecated. Use "
385  "exclusive_scan_over_group instead.")
386 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
387  detail::is_vector_arithmetic<T>::value &&
388  detail::is_native_op<T, BinaryOperation>::value),
389  T> exclusive_scan(Group g, T x, BinaryOperation binary_op) {
390  return exclusive_scan_over_group(g, x, binary_op);
391 }
392 
393 template <typename Group, typename V, typename T, class BinaryOperation>
394 __SYCL2020_DEPRECATED("ext::oneapi::exclusive_scan is deprecated. Use "
395  "exclusive_scan_over_group instead.")
396 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
397  detail::is_vector_arithmetic<V>::value &&
398  detail::is_vector_arithmetic<T>::value &&
399  detail::is_native_op<V, BinaryOperation>::value &&
400  detail::is_native_op<T, BinaryOperation>::value),
401  T> exclusive_scan(Group g, V x, T init,
402  BinaryOperation binary_op) {
403  return exclusive_scan_over_group(g, x, init, binary_op);
404 }
405 
406 template <typename Group, typename V, typename T, class BinaryOperation>
407 __SYCL2020_DEPRECATED("ext::oneapi::exclusive_scan is deprecated. Use "
408  "exclusive_scan_over_group instead.")
409 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
410  detail::is_scalar_arithmetic<V>::value &&
411  detail::is_scalar_arithmetic<T>::value &&
412  detail::is_native_op<V, BinaryOperation>::value &&
413  detail::is_native_op<T, BinaryOperation>::value),
414  T> exclusive_scan(Group g, V x, T init,
415  BinaryOperation binary_op) {
416  return exclusive_scan_over_group(g, x, init, binary_op);
417 }
418 
419 template <typename Group, typename InPtr, typename OutPtr, typename T,
420  class BinaryOperation>
421 __SYCL2020_DEPRECATED("ext::oneapi::exclusive_scan is deprecated. Use "
422  "joint_exclusive_scan instead.")
423 detail::enable_if_t<
424  (detail::is_generic_group<Group>::value &&
425  detail::is_pointer<InPtr>::value && detail::is_pointer<OutPtr>::value &&
426  detail::is_arithmetic<
427  typename detail::remove_pointer<InPtr>::type>::value &&
428  detail::is_arithmetic<T>::value &&
429  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
430  BinaryOperation>::value &&
431  detail::is_native_op<T, BinaryOperation>::value),
432  OutPtr> exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
433  T init, BinaryOperation binary_op) {
434  return joint_exclusive_scan(g, first, last, result, init, binary_op);
435 }
436 
437 template <typename Group, typename InPtr, typename OutPtr,
438  class BinaryOperation>
439 __SYCL2020_DEPRECATED("ext::oneapi::exclusive_scan is deprecated. Use "
440  "joint_exclusive_scan instead.")
441 detail::enable_if_t<
442  (detail::is_generic_group<Group>::value &&
443  detail::is_pointer<InPtr>::value && detail::is_pointer<OutPtr>::value &&
444  detail::is_arithmetic<
445  typename detail::remove_pointer<InPtr>::type>::value &&
446  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
447  BinaryOperation>::value),
448  OutPtr> exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
449  BinaryOperation binary_op) {
450  return joint_exclusive_scan(g, first, last, result, binary_op);
451 }
452 
453 template <typename Group, typename T, class BinaryOperation>
454 __SYCL2020_DEPRECATED("ext::oneapi::inclusive_scan is deprecated. Use "
455  "inclusive_scan_over_group instead.")
456 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
457  detail::is_vector_arithmetic<T>::value &&
458  detail::is_native_op<T, BinaryOperation>::value),
459  T> inclusive_scan(Group g, T x, BinaryOperation binary_op) {
460  return inclusive_scan_over_group(g, x, binary_op);
461 }
462 
463 template <typename Group, typename T, class BinaryOperation>
464 __SYCL2020_DEPRECATED("ext::oneapi::inclusive_scan is deprecated. Use "
465  "inclusive_scan_over_group instead.")
466 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
467  detail::is_scalar_arithmetic<T>::value &&
468  detail::is_native_op<T, BinaryOperation>::value),
469  T> inclusive_scan(Group g, T x, BinaryOperation binary_op) {
470  return inclusive_scan_over_group(g, x, binary_op);
471 }
472 
473 template <typename Group, typename V, class BinaryOperation, typename T>
474 __SYCL2020_DEPRECATED("ext::oneapi::inclusive_scan is deprecated. Use "
475  "inclusive_scan_over_group instead.")
476 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
477  detail::is_scalar_arithmetic<V>::value &&
478  detail::is_scalar_arithmetic<T>::value &&
479  detail::is_native_op<V, BinaryOperation>::value &&
480  detail::is_native_op<T, BinaryOperation>::value),
481  T> inclusive_scan(Group g, V x, BinaryOperation binary_op,
482  T init) {
483  return inclusive_scan_over_group(g, x, binary_op, init);
484 }
485 
486 template <typename Group, typename V, class BinaryOperation, typename T>
487 __SYCL2020_DEPRECATED("ext::oneapi::inclusive_scan is deprecated. Use "
488  "inclusive_scan_over_group instead.")
489 detail::enable_if_t<(detail::is_generic_group<Group>::value &&
490  detail::is_vector_arithmetic<V>::value &&
491  detail::is_vector_arithmetic<T>::value &&
492  detail::is_native_op<V, BinaryOperation>::value &&
493  detail::is_native_op<T, BinaryOperation>::value),
494  T> inclusive_scan(Group g, V x, BinaryOperation binary_op,
495  T init) {
496  return inclusive_scan_over_group(g, x, binary_op, init);
497 }
498 
499 template <typename Group, typename InPtr, typename OutPtr,
500  class BinaryOperation, typename T>
501 __SYCL2020_DEPRECATED("ext::oneapi::inclusive_scan is deprecated. Use "
502  "joint_inclusive_scan instead.")
503 detail::enable_if_t<
504  (detail::is_generic_group<Group>::value &&
505  detail::is_pointer<InPtr>::value && detail::is_pointer<OutPtr>::value &&
506  detail::is_arithmetic<
507  typename detail::remove_pointer<InPtr>::type>::value &&
508  detail::is_arithmetic<T>::value &&
509  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
510  BinaryOperation>::value &&
511  detail::is_native_op<T, BinaryOperation>::value),
512  OutPtr> inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
513  BinaryOperation binary_op, T init) {
514  return joint_inclusive_scan(g, first, last, result, binary_op, init);
515 }
516 
517 template <typename Group, typename InPtr, typename OutPtr,
518  class BinaryOperation>
519 __SYCL2020_DEPRECATED("ext::oneapi::inclusive_scan is deprecated. Use "
520  "joint_inclusive_scan instead.")
521 detail::enable_if_t<
522  (detail::is_generic_group<Group>::value &&
523  detail::is_pointer<InPtr>::value && detail::is_pointer<OutPtr>::value &&
524  detail::is_arithmetic<
525  typename detail::remove_pointer<InPtr>::type>::value &&
526  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
527  BinaryOperation>::value),
528  OutPtr> inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
529  BinaryOperation binary_op) {
530  return joint_inclusive_scan(g, first, last, result, binary_op);
531 }
532 
533 template <typename Group>
535  "ext::oneapi::leader free function is deprecated. Use member function "
536  "leader of the sycl::group/sycl::sub_group instead.")
537 detail::enable_if_t<detail::is_generic_group<Group>::value, bool> leader(
538  Group g) {
539 #ifdef __SYCL_DEVICE_ONLY__
540  typename Group::linear_id_type linear_id =
542  return (linear_id == 0);
543 #else
544  (void)g;
545  throw runtime_error("Group algorithms are not supported on host device.",
546  PI_ERROR_INVALID_DEVICE);
547 #endif
548 }
549 
550 } // namespace ext::oneapi
551 
552 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
553 } // namespace sycl
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL2020_DEPRECATED(message)
id< 3 > linear_id_to_id(range< 3 > r, size_t linear_id)
typename std::enable_if< B, T >::type enable_if_t
Group::linear_id_type get_local_linear_id(Group g)
sycl::detail::enable_if_t< sycl::detail::is_scalar_arithmetic< T >::value, T > EnableIfIsScalarArithmetic
sycl::detail::enable_if_t< sycl::detail::is_vector_arithmetic< T >::value, T > EnableIfIsVectorArithmetic
sycl::detail::enable_if_t< sycl::detail::is_vector_arithmetic< T >::value &&sycl::detail::is_native_op< T, BinaryOperation >::value, T > EnableIfIsVectorArithmeticNativeOp
sycl::detail::enable_if_t< std::is_trivially_copyable< T >::value &&!sycl::detail::is_vector_arithmetic< T >::value, T > EnableIfIsTriviallyCopyable
sycl::detail::enable_if_t< sycl::detail::is_scalar_arithmetic< T >::value &&sycl::detail::is_native_op< T, BinaryOperation >::value, T > EnableIfIsScalarArithmeticNativeOp
sycl::detail::enable_if_t<(!sycl::detail::is_scalar_arithmetic< T >::value &&!sycl::detail::is_vector_arithmetic< T >::value &&std::is_trivially_copyable< T >::value)||!sycl::detail::is_native_op< T, BinaryOperation >::value, T > EnableIfIsNonNativeOp
sycl::detail::enable_if_t< sycl::detail::is_pointer< Ptr >::value, T > EnableIfIsPointer
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer< Ptr >::value &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< Ptr >::type >::value &&detail::is_arithmetic_or_complex< T >::value &&detail::is_native_op< typename detail::remove_pointer< Ptr >::type, BinaryOperation >::value &&detail::is_plus_if_complex< typename detail::remove_pointer< Ptr >::type, BinaryOperation >::value &&detail::is_plus_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)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&(detail::is_scalar_arithmetic< V >::value||detail::is_complex< V >::value) &&(detail::is_scalar_arithmetic< T >::value||detail::is_complex< T >::value) &&detail::is_native_op< V, BinaryOperation >::value &&detail::is_native_op< T, BinaryOperation >::value &&detail::is_plus_if_complex< V, BinaryOperation >::value &&detail::is_plus_if_complex< T, BinaryOperation >::value), T > exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op)
detail::enable_if_t< is_group_v< std::decay_t< Group > >, bool > none_of_group(Group g, T x, Predicate pred)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer< Ptr >::value), bool > joint_all_of(Group g, Ptr first, Ptr last, Predicate pred)
detail::enable_if_t< is_group_v< Group >, bool > any_of_group(Group g, T x, Predicate pred)
detail::enable_if_t< is_group_v< std::decay_t< Group > >, bool > all_of_group(Group g, T x, Predicate pred)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer< Ptr >::value), bool > joint_any_of(Group g, Ptr first, Ptr last, Predicate pred)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer< Ptr >::value), bool > joint_none_of(Group g, Ptr first, Ptr last, Predicate pred)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_vector_arithmetic< V >::value &&detail::is_vector_arithmetic< T >::value &&detail::is_native_op< V, BinaryOperation >::value &&detail::is_native_op< T, BinaryOperation >::value), T > reduce_over_group(Group g, V x, T init, BinaryOperation binary_op)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_vector_arithmetic< V >::value &&detail::is_vector_arithmetic< T >::value &&detail::is_native_op< V, BinaryOperation >::value &&detail::is_native_op< T, BinaryOperation >::value), T > inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer< InPtr >::value &&detail::is_pointer< OutPtr >::value &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< InPtr >::type >::value &&detail::is_native_op< typename detail::remove_pointer< InPtr >::type, BinaryOperation >::value &&detail::is_plus_if_complex< typename detail::remove_pointer< InPtr >::type, BinaryOperation >::value), OutPtr > joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op)
detail::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer< InPtr >::value &&detail::is_pointer< OutPtr >::value &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< InPtr >::type >::value &&detail::is_native_op< typename detail::remove_pointer< InPtr >::type, BinaryOperation >::value &&detail::is_plus_if_complex< typename detail::remove_pointer< InPtr >::type, BinaryOperation >::value), OutPtr > joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op)
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept
bool none_of(const simd_mask< _Tp, _Abi > &) noexcept
_Tp reduce(const simd< _Tp, _Abi > &, _BinaryOp=_BinaryOp())