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