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 #include <complex>
11 
12 #include <CL/__spirv/spirv_ops.hpp>
15 #include <sycl/builtins.hpp>
16 #include <sycl/detail/spirv.hpp>
19 #include <sycl/functional.hpp>
20 #include <sycl/group.hpp>
21 #include <sycl/group_barrier.hpp>
22 #include <sycl/known_identity.hpp>
23 #include <sycl/nd_item.hpp>
24 #include <sycl/sub_group.hpp>
25 
26 namespace sycl {
28 namespace detail {
29 
30 // ---- linear_id_to_id
31 template <int Dimensions>
33 template <> inline id<1> linear_id_to_id(range<1>, size_t linear_id) {
34  return id<1>(linear_id);
35 }
36 template <> inline id<2> linear_id_to_id(range<2> r, size_t linear_id) {
37  id<2> result;
38  result[0] = linear_id / r[1];
39  result[1] = linear_id % r[1];
40  return result;
41 }
42 template <> inline id<3> linear_id_to_id(range<3> r, size_t linear_id) {
43  id<3> result;
44  result[0] = linear_id / (r[1] * r[2]);
45  result[1] = (linear_id % (r[1] * r[2])) / r[2];
46  result[2] = linear_id % r[2];
47  return result;
48 }
49 
50 // ---- get_local_linear_range
51 template <typename Group> size_t get_local_linear_range(Group g);
52 template <> inline size_t get_local_linear_range<group<1>>(group<1> g) {
53  return g.get_local_range(0);
54 }
55 template <> inline size_t get_local_linear_range<group<2>>(group<2> g) {
56  return g.get_local_range(0) * g.get_local_range(1);
57 }
58 template <> inline size_t get_local_linear_range<group<3>>(group<3> g) {
59  return g.get_local_range(0) * g.get_local_range(1) * g.get_local_range(2);
60 }
61 template <>
62 inline size_t
63 get_local_linear_range<ext::oneapi::sub_group>(ext::oneapi::sub_group g) {
64  return g.get_local_range()[0];
65 }
66 
67 // ---- get_local_linear_id
68 template <typename Group>
69 inline typename Group::linear_id_type get_local_linear_id(Group g);
70 
71 #ifdef __SYCL_DEVICE_ONLY__
72 #define __SYCL_GROUP_GET_LOCAL_LINEAR_ID(D) \
73  template <> \
74  inline group<D>::linear_id_type get_local_linear_id<group<D>>(group<D>) { \
75  nd_item<D> it = sycl::detail::Builder::getNDItem<D>(); \
76  return it.get_local_linear_id(); \
77  }
78 __SYCL_GROUP_GET_LOCAL_LINEAR_ID(1);
79 __SYCL_GROUP_GET_LOCAL_LINEAR_ID(2);
80 __SYCL_GROUP_GET_LOCAL_LINEAR_ID(3);
81 #undef __SYCL_GROUP_GET_LOCAL_LINEAR_ID
82 #endif // __SYCL_DEVICE_ONLY__
83 
84 template <>
86 get_local_linear_id<ext::oneapi::sub_group>(ext::oneapi::sub_group g) {
87  return g.get_local_id()[0];
88 }
89 
90 // ---- is_native_op
91 template <typename T>
93  type_list<sycl::plus<T>, sycl::bit_or<T>, sycl::bit_xor<T>,
94  sycl::bit_and<T>, sycl::maximum<T>, sycl::minimum<T>,
95  sycl::multiplies<T>>;
96 
97 template <typename T, typename BinaryOperation> struct is_native_op {
98  static constexpr bool value =
101 };
102 
103 // ---- is_plus
104 template <typename T, typename BinaryOperation>
105 using is_plus = std::integral_constant<
106  bool, std::is_same<BinaryOperation, sycl::plus<T>>::value ||
107  std::is_same<BinaryOperation, sycl::plus<void>>::value>;
108 
109 // ---- is_complex
110 // NOTE: std::complex<long double> not yet supported by group algorithms.
111 template <typename T>
113  : std::integral_constant<bool,
114  std::is_same<T, std::complex<float>>::value ||
115  std::is_same<T, std::complex<double>>::value> {
116 };
117 
118 // ---- is_arithmetic_or_complex
119 template <typename T>
121  std::integral_constant<bool, sycl::detail::is_complex<T>::value ||
122  sycl::detail::is_arithmetic<T>::value>;
123 // ---- is_plus_if_complex
124 template <typename T, typename BinaryOperation>
126  std::integral_constant<bool, (is_complex<T>::value
128  : std::true_type::value)>;
129 
130 // ---- identity_for_ga_op
131 // the group algorithms support std::complex, limited to sycl::plus operation
132 // get the correct identity for group algorithm operation.
133 // TODO: identiy_for_ga_op should be replaced with known_identity once the other
134 // callers of known_identity support complex numbers.
135 template <typename T, class BinaryOperation>
136 constexpr detail::enable_if_t<
139  return {0, 0};
140 }
141 
142 template <typename T, class BinaryOperation>
144  return sycl::known_identity_v<BinaryOperation, T>;
145 }
146 
147 // ---- for_each
148 template <typename Group, typename Ptr, class Function>
149 Function for_each(Group g, Ptr first, Ptr last, Function f) {
150 #ifdef __SYCL_DEVICE_ONLY__
151  ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
152  ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
153  for (Ptr p = first + offset; p < last; p += stride) {
154  f(*p);
155  }
156  return f;
157 #else
158  (void)g;
159  (void)first;
160  (void)last;
161  (void)f;
162  throw runtime_error("Group algorithms are not supported on host device.",
163  PI_ERROR_INVALID_DEVICE);
164 #endif
165 }
166 } // namespace detail
167 
168 // ---- reduce_over_group
169 // three argument variant is specialized thrice:
170 // scalar arithmetic, complex (plus only), and vector arithmetic
171 
172 template <typename Group, typename T, class BinaryOperation>
173 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
174  detail::is_scalar_arithmetic<T>::value &&
175  detail::is_native_op<T, BinaryOperation>::value),
176  T>
177 reduce_over_group(Group, T x, BinaryOperation binary_op) {
178  // FIXME: Do not special-case for half precision
179  static_assert(
180  std::is_same<decltype(binary_op(x, x)), T>::value ||
181  (std::is_same<T, half>::value &&
182  std::is_same<decltype(binary_op(x, x)), float>::value),
183  "Result type of binary_op must match reduction accumulation type.");
184 #ifdef __SYCL_DEVICE_ONLY__
185  return sycl::detail::calc<T, __spv::GroupOperation::Reduce,
186  sycl::detail::spirv::group_scope<Group>::value>(
187  typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
188 #else
189  throw runtime_error("Group algorithms are not supported on host device.",
190  PI_ERROR_INVALID_DEVICE);
191 #endif
192 }
193 
194 // complex specialization. T is std::complex<float> or similar.
195 // binary op is sycl::plus<std::complex<float>>
196 template <typename Group, typename T, class BinaryOperation>
197 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
198  detail::is_complex<T>::value &&
199  detail::is_native_op<T, sycl::plus<T>>::value &&
200  detail::is_plus<T, BinaryOperation>::value),
201  T>
202 reduce_over_group(Group g, T x, BinaryOperation binary_op) {
203 #ifdef __SYCL_DEVICE_ONLY__
204  T result;
205  result.real(reduce_over_group(g, x.real(), sycl::plus<>()));
206  result.imag(reduce_over_group(g, x.imag(), sycl::plus<>()));
207  return result;
208 #else
209  (void)g;
210  (void)x;
211  (void)binary_op;
212  throw runtime_error("Group algorithms are not supported on host device.",
213  PI_ERROR_INVALID_DEVICE);
214 #endif
215 }
216 
217 template <typename Group, typename T, class BinaryOperation>
218 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
219  detail::is_vector_arithmetic<T>::value &&
220  detail::is_native_op<T, BinaryOperation>::value),
221  T>
222 reduce_over_group(Group g, T x, BinaryOperation binary_op) {
223  // FIXME: Do not special-case for half precision
224  static_assert(
225  std::is_same<decltype(binary_op(x[0], x[0])),
226  typename T::element_type>::value ||
227  (std::is_same<T, half>::value &&
228  std::is_same<decltype(binary_op(x[0], x[0])), float>::value),
229  "Result type of binary_op must match reduction accumulation type.");
230  T result;
231  for (int s = 0; s < x.size(); ++s) {
232  result[s] = reduce_over_group(g, x[s], binary_op);
233  }
234  return result;
235 }
236 
237 // four argument variant of reduce_over_group specialized twice
238 // (scalar arithmetic || complex), and vector_arithmetic
239 template <typename Group, typename V, typename T, class BinaryOperation>
241  (is_group_v<std::decay_t<Group>> &&
242  (detail::is_scalar_arithmetic<V>::value || detail::is_complex<V>::value) &&
243  (detail::is_scalar_arithmetic<T>::value || detail::is_complex<T>::value) &&
244  detail::is_native_op<V, BinaryOperation>::value &&
245  detail::is_native_op<T, BinaryOperation>::value &&
246  detail::is_plus_if_complex<T, BinaryOperation>::value &&
247  detail::is_plus_if_complex<V, BinaryOperation>::value),
248  T>
249 reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
250  // FIXME: Do not special-case for half precision
251  static_assert(
252  std::is_same<decltype(binary_op(init, x)), T>::value ||
253  (std::is_same<T, half>::value &&
254  std::is_same<decltype(binary_op(init, x)), float>::value),
255  "Result type of binary_op must match reduction accumulation type.");
256 #ifdef __SYCL_DEVICE_ONLY__
257  return binary_op(init, reduce_over_group(g, x, binary_op));
258 #else
259  (void)g;
260  throw runtime_error("Group algorithms are not supported on host device.",
261  PI_ERROR_INVALID_DEVICE);
262 #endif
263 }
264 
265 template <typename Group, typename V, typename T, class BinaryOperation>
266 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
267  detail::is_vector_arithmetic<V>::value &&
268  detail::is_vector_arithmetic<T>::value &&
269  detail::is_native_op<V, BinaryOperation>::value &&
270  detail::is_native_op<T, BinaryOperation>::value),
271  T>
272 reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
273  // FIXME: Do not special-case for half precision
274  static_assert(
275  std::is_same<decltype(binary_op(init[0], x[0])),
276  typename T::element_type>::value ||
277  (std::is_same<T, half>::value &&
278  std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
279  "Result type of binary_op must match reduction accumulation type.");
280 #ifdef __SYCL_DEVICE_ONLY__
281  T result = init;
282  for (int s = 0; s < x.size(); ++s) {
283  result[s] = binary_op(init[s], reduce_over_group(g, x[s], binary_op));
284  }
285  return result;
286 #else
287  (void)g;
288  throw runtime_error("Group algorithms are not supported on host device.",
289  PI_ERROR_INVALID_DEVICE);
290 #endif
291 }
292 
293 // ---- joint_reduce
294 template <typename Group, typename Ptr, class BinaryOperation>
296  (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value &&
298  typename detail::remove_pointer<Ptr>::type>::value &&
299  detail::is_plus_if_complex<typename detail::remove_pointer<Ptr>::type,
300  BinaryOperation>::value),
301  typename detail::remove_pointer<Ptr>::type>
302 joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
303 #ifdef __SYCL_DEVICE_ONLY__
304  using T = typename detail::remove_pointer<Ptr>::type;
305  T init = detail::identity_for_ga_op<T, BinaryOperation>();
306  return joint_reduce(g, first, last, init, binary_op);
307 #else
308  (void)g;
309  (void)first;
310  (void)last;
311  (void)binary_op;
312  throw runtime_error("Group algorithms are not supported on host device.",
313  PI_ERROR_INVALID_DEVICE);
314 #endif
315 }
316 
317 template <typename Group, typename Ptr, typename T, class BinaryOperation>
319  (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value &&
321  typename detail::remove_pointer<Ptr>::type>::value &&
322  detail::is_arithmetic_or_complex<T>::value &&
323  detail::is_native_op<typename detail::remove_pointer<Ptr>::type,
324  BinaryOperation>::value &&
325  detail::is_plus_if_complex<typename detail::remove_pointer<Ptr>::type,
326  BinaryOperation>::value &&
327  detail::is_plus_if_complex<T, BinaryOperation>::value &&
328  detail::is_native_op<T, BinaryOperation>::value),
329  T>
330 joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
331  // FIXME: Do not special-case for half precision
332  static_assert(
333  std::is_same<decltype(binary_op(init, *first)), T>::value ||
334  (std::is_same<T, half>::value &&
335  std::is_same<decltype(binary_op(init, *first)), float>::value),
336  "Result type of binary_op must match reduction accumulation type.");
337 #ifdef __SYCL_DEVICE_ONLY__
338  T partial = detail::identity_for_ga_op<T, BinaryOperation>();
340  g, first, last, [&](const typename detail::remove_pointer<Ptr>::type &x) {
341  partial = binary_op(partial, x);
342  });
343  return reduce_over_group(g, partial, init, binary_op);
344 #else
345  (void)g;
346  (void)last;
347  throw runtime_error("Group algorithms are not supported on host device.",
348  PI_ERROR_INVALID_DEVICE);
349 #endif
350 }
351 
352 // ---- any_of_group
353 template <typename Group>
354 detail::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
355 any_of_group(Group, bool pred) {
356 #ifdef __SYCL_DEVICE_ONLY__
357  return sycl::detail::spirv::GroupAny<Group>(pred);
358 #else
359  (void)pred;
360  throw runtime_error("Group algorithms are not supported on host device.",
361  PI_ERROR_INVALID_DEVICE);
362 #endif
363 }
364 
365 template <typename Group, typename T, class Predicate>
367  Predicate pred) {
368  return any_of_group(g, pred(x));
369 }
370 
371 // ---- joint_any_of
372 template <typename Group, typename Ptr, class Predicate>
374  (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value), bool>
375 joint_any_of(Group g, Ptr first, Ptr last, Predicate pred) {
376 #ifdef __SYCL_DEVICE_ONLY__
377  using T = typename detail::remove_pointer<Ptr>::type;
378  bool partial = false;
379  sycl::detail::for_each(g, first, last, [&](T &x) { partial |= pred(x); });
380  return any_of_group(g, partial);
381 #else
382  (void)g;
383  (void)first;
384  (void)last;
385  (void)pred;
386  throw runtime_error("Group algorithms are not supported on host device.",
387  PI_ERROR_INVALID_DEVICE);
388 #endif
389 }
390 
391 // ---- all_of_group
392 template <typename Group>
393 detail::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
394 all_of_group(Group, bool pred) {
395 #ifdef __SYCL_DEVICE_ONLY__
396  return sycl::detail::spirv::GroupAll<Group>(pred);
397 #else
398  (void)pred;
399  throw runtime_error("Group algorithms are not supported on host device.",
400  PI_ERROR_INVALID_DEVICE);
401 #endif
402 }
403 
404 template <typename Group, typename T, class Predicate>
405 detail::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
406 all_of_group(Group g, T x, Predicate pred) {
407  return all_of_group(g, pred(x));
408 }
409 
410 // ---- joint_all_of
411 template <typename Group, typename Ptr, class Predicate>
413  (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value), bool>
414 joint_all_of(Group g, Ptr first, Ptr last, Predicate pred) {
415 #ifdef __SYCL_DEVICE_ONLY__
416  using T = typename detail::remove_pointer<Ptr>::type;
417  bool partial = true;
418  sycl::detail::for_each(g, first, last, [&](T &x) { partial &= pred(x); });
419  return all_of_group(g, partial);
420 #else
421  (void)g;
422  (void)first;
423  (void)last;
424  (void)pred;
425  throw runtime_error("Group algorithms are not supported on host device.",
426  PI_ERROR_INVALID_DEVICE);
427 #endif
428 }
429 
430 // ---- none_of_group
431 template <typename Group>
432 detail::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
433 none_of_group(Group, bool pred) {
434 #ifdef __SYCL_DEVICE_ONLY__
435  return sycl::detail::spirv::GroupAll<Group>(!pred);
436 #else
437  (void)pred;
438  throw runtime_error("Group algorithms are not supported on host device.",
439  PI_ERROR_INVALID_DEVICE);
440 #endif
441 }
442 
443 template <typename Group, typename T, class Predicate>
444 detail::enable_if_t<is_group_v<std::decay_t<Group>>, bool>
445 none_of_group(Group g, T x, Predicate pred) {
446  return none_of_group(g, pred(x));
447 }
448 
449 // ---- joint_none_of
450 template <typename Group, typename Ptr, class Predicate>
452  (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value), bool>
453 joint_none_of(Group g, Ptr first, Ptr last, Predicate pred) {
454 #ifdef __SYCL_DEVICE_ONLY__
455  return !joint_any_of(g, first, last, pred);
456 #else
457  (void)g;
458  (void)first;
459  (void)last;
460  (void)pred;
461  throw runtime_error("Group algorithms are not supported on host device.",
462  PI_ERROR_INVALID_DEVICE);
463 #endif
464 }
465 
466 // ---- shift_group_left
467 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
468 // copyable.
469 template <typename Group, typename T>
470 detail::enable_if_t<(std::is_same<std::decay_t<Group>, sub_group>::value &&
471  (std::is_trivially_copyable<T>::value ||
472  detail::is_vec<T>::value)),
473  T>
474 shift_group_left(Group, T x, typename Group::linear_id_type delta = 1) {
475 #ifdef __SYCL_DEVICE_ONLY__
476  return sycl::detail::spirv::SubgroupShuffleDown(x, delta);
477 #else
478  (void)x;
479  (void)delta;
480  throw runtime_error("Sub-groups are not supported on host device.",
481  PI_ERROR_INVALID_DEVICE);
482 #endif
483 }
484 
485 // ---- shift_group_right
486 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
487 // copyable.
488 template <typename Group, typename T>
489 detail::enable_if_t<(std::is_same<std::decay_t<Group>, sub_group>::value &&
490  (std::is_trivially_copyable<T>::value ||
491  detail::is_vec<T>::value)),
492  T>
493 shift_group_right(Group, T x, typename Group::linear_id_type delta = 1) {
494 #ifdef __SYCL_DEVICE_ONLY__
495  return sycl::detail::spirv::SubgroupShuffleUp(x, delta);
496 #else
497  (void)x;
498  (void)delta;
499  throw runtime_error("Sub-groups are not supported on host device.",
500  PI_ERROR_INVALID_DEVICE);
501 #endif
502 }
503 
504 // ---- permute_group_by_xor
505 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
506 // copyable.
507 template <typename Group, typename T>
508 detail::enable_if_t<(std::is_same<std::decay_t<Group>, sub_group>::value &&
509  (std::is_trivially_copyable<T>::value ||
510  detail::is_vec<T>::value)),
511  T>
512 permute_group_by_xor(Group, T x, typename Group::linear_id_type mask) {
513 #ifdef __SYCL_DEVICE_ONLY__
514  return sycl::detail::spirv::SubgroupShuffleXor(x, mask);
515 #else
516  (void)x;
517  (void)mask;
518  throw runtime_error("Sub-groups are not supported on host device.",
519  PI_ERROR_INVALID_DEVICE);
520 #endif
521 }
522 
523 // ---- select_from_group
524 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
525 // copyable.
526 template <typename Group, typename T>
527 detail::enable_if_t<(std::is_same<std::decay_t<Group>, sub_group>::value &&
528  (std::is_trivially_copyable<T>::value ||
529  detail::is_vec<T>::value)),
530  T>
531 select_from_group(Group, T x, typename Group::id_type local_id) {
532 #ifdef __SYCL_DEVICE_ONLY__
533  return sycl::detail::spirv::SubgroupShuffle(x, local_id);
534 #else
535  (void)x;
536  (void)local_id;
537  throw runtime_error("Sub-groups are not supported on host device.",
538  PI_ERROR_INVALID_DEVICE);
539 #endif
540 }
541 
542 // ---- group_broadcast
543 // TODO: remove check for detail::is_vec<T> once sycl::vec is trivially
544 // copyable.
545 template <typename Group, typename T>
546 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
547  (std::is_trivially_copyable<T>::value ||
548  detail::is_vec<T>::value)),
549  T>
550 group_broadcast(Group, T x, typename Group::id_type local_id) {
551 #ifdef __SYCL_DEVICE_ONLY__
552  return sycl::detail::spirv::GroupBroadcast<Group>(x, local_id);
553 #else
554  (void)x;
555  (void)local_id;
556  throw runtime_error("Group algorithms are not supported on host device.",
557  PI_ERROR_INVALID_DEVICE);
558 #endif
559 }
560 
561 template <typename Group, typename T>
562 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
563  (std::is_trivially_copyable<T>::value ||
564  detail::is_vec<T>::value)),
565  T>
566 group_broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
567 #ifdef __SYCL_DEVICE_ONLY__
568  return group_broadcast(
569  g, x,
570  sycl::detail::linear_id_to_id(g.get_local_range(), linear_local_id));
571 #else
572  (void)g;
573  (void)x;
574  (void)linear_local_id;
575  throw runtime_error("Group algorithms are not supported on host device.",
576  PI_ERROR_INVALID_DEVICE);
577 #endif
578 }
579 
580 template <typename Group, typename T>
581 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
582  (std::is_trivially_copyable<T>::value ||
583  detail::is_vec<T>::value)),
584  T>
585 group_broadcast(Group g, T x) {
586 #ifdef __SYCL_DEVICE_ONLY__
587  return group_broadcast(g, x, 0);
588 #else
589  (void)g;
590  (void)x;
591  throw runtime_error("Group algorithms are not supported on host device.",
592  PI_ERROR_INVALID_DEVICE);
593 #endif
594 }
595 
596 // ---- exclusive_scan_over_group
597 // this function has two overloads, one with three arguments and one with four
598 // the three argument version is specialized thrice: scalar, complex, and
599 // vector
600 template <typename Group, typename T, class BinaryOperation>
601 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
602  detail::is_scalar_arithmetic<T>::value &&
603  detail::is_native_op<T, BinaryOperation>::value),
604  T>
605 exclusive_scan_over_group(Group, T x, BinaryOperation binary_op) {
606  // FIXME: Do not special-case for half precision
607  static_assert(std::is_same<decltype(binary_op(x, x)), T>::value ||
608  (std::is_same<T, half>::value &&
609  std::is_same<decltype(binary_op(x, x)), float>::value),
610  "Result type of binary_op must match scan accumulation type.");
611 #ifdef __SYCL_DEVICE_ONLY__
612  return sycl::detail::calc<T, __spv::GroupOperation::ExclusiveScan,
613  sycl::detail::spirv::group_scope<Group>::value>(
614  typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
615 #else
616  throw runtime_error("Group algorithms are not supported on host device.",
617  PI_ERROR_INVALID_DEVICE);
618 #endif
619 }
620 
621 // complex specialization. T is std::complex<float> or similar.
622 // binary op is sycl::plus<std::complex<float>>
623 template <typename Group, typename T, class BinaryOperation>
624 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
625  detail::is_complex<T>::value &&
626  detail::is_native_op<T, sycl::plus<T>>::value &&
627  detail::is_plus<T, BinaryOperation>::value),
628  T>
629 exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
630 #ifdef __SYCL_DEVICE_ONLY__
631  T result;
632  result.real(exclusive_scan_over_group(g, x.real(), sycl::plus<>()));
633  result.imag(exclusive_scan_over_group(g, x.imag(), sycl::plus<>()));
634  return result;
635 #else
636  (void)g;
637  (void)x;
638  (void)binary_op;
639  throw runtime_error("Group algorithms are not supported on host device.",
640  PI_ERROR_INVALID_DEVICE);
641 #endif
642 }
643 
644 template <typename Group, typename T, class BinaryOperation>
645 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
646  detail::is_vector_arithmetic<T>::value &&
647  detail::is_native_op<T, BinaryOperation>::value),
648  T>
649 exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
650  // FIXME: Do not special-case for half precision
651  static_assert(
652  std::is_same<decltype(binary_op(x[0], x[0])),
653  typename T::element_type>::value ||
654  (std::is_same<T, half>::value &&
655  std::is_same<decltype(binary_op(x[0], x[0])), float>::value),
656  "Result type of binary_op must match scan accumulation type.");
657  T result;
658  for (int s = 0; s < x.size(); ++s) {
659  result[s] = exclusive_scan_over_group(g, x[s], binary_op);
660  }
661  return result;
662 }
663 
664 // four argument version of exclusive_scan_over_group is specialized twice
665 // once for vector_arithmetic, once for (scalar_arithmetic || complex)
666 template <typename Group, typename V, typename T, class BinaryOperation>
667 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
668  detail::is_vector_arithmetic<V>::value &&
669  detail::is_vector_arithmetic<T>::value &&
670  detail::is_native_op<V, BinaryOperation>::value &&
671  detail::is_native_op<T, BinaryOperation>::value),
672  T>
673 exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
674  // FIXME: Do not special-case for half precision
675  static_assert(
676  std::is_same<decltype(binary_op(init[0], x[0])),
677  typename T::element_type>::value ||
678  (std::is_same<T, half>::value &&
679  std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
680  "Result type of binary_op must match scan accumulation type.");
681  T result;
682  for (int s = 0; s < x.size(); ++s) {
683  result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op);
684  }
685  return result;
686 }
687 
688 template <typename Group, typename V, typename T, class BinaryOperation>
690  (is_group_v<std::decay_t<Group>> &&
691  (detail::is_scalar_arithmetic<V>::value || detail::is_complex<V>::value) &&
692  (detail::is_scalar_arithmetic<T>::value || detail::is_complex<T>::value) &&
693  detail::is_native_op<V, BinaryOperation>::value &&
694  detail::is_native_op<T, BinaryOperation>::value &&
695  detail::is_plus_if_complex<V, BinaryOperation>::value &&
696  detail::is_plus_if_complex<T, BinaryOperation>::value),
697  T>
698 exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
699  // FIXME: Do not special-case for half precision
700  static_assert(std::is_same<decltype(binary_op(init, x)), T>::value ||
701  (std::is_same<T, half>::value &&
702  std::is_same<decltype(binary_op(init, x)), float>::value),
703  "Result type of binary_op must match scan accumulation type.");
704 #ifdef __SYCL_DEVICE_ONLY__
705  typename Group::linear_id_type local_linear_id =
707  if (local_linear_id == 0) {
708  x = binary_op(init, x);
709  }
710  T scan = exclusive_scan_over_group(g, x, binary_op);
711  if (local_linear_id == 0) {
712  scan = init;
713  }
714  return scan;
715 #else
716  (void)g;
717  throw runtime_error("Group algorithms are not supported on host device.",
718  PI_ERROR_INVALID_DEVICE);
719 #endif
720 }
721 
722 // ---- joint_exclusive_scan
723 template <typename Group, typename InPtr, typename OutPtr, typename T,
724  class BinaryOperation>
726  (is_group_v<std::decay_t<Group>> && detail::is_pointer<InPtr>::value &&
727  detail::is_pointer<OutPtr>::value &&
729  typename detail::remove_pointer<InPtr>::type>::value &&
730  detail::is_arithmetic_or_complex<T>::value &&
731  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
732  BinaryOperation>::value &&
733  detail::is_native_op<T, BinaryOperation>::value &&
734  detail::is_plus_if_complex<typename detail::remove_pointer<InPtr>::type,
735  BinaryOperation>::value &&
736  detail::is_plus_if_complex<T, BinaryOperation>::value),
737  OutPtr>
738 joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
739  BinaryOperation binary_op) {
740  // FIXME: Do not special-case for half precision
741  static_assert(
742  std::is_same<decltype(binary_op(*first, *first)), T>::value ||
743  (std::is_same<T, half>::value &&
744  std::is_same<decltype(binary_op(*first, *first)), float>::value),
745  "Result type of binary_op must match scan accumulation type.");
746 #ifdef __SYCL_DEVICE_ONLY__
747  ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
748  ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
749  ptrdiff_t N = last - first;
750  auto roundup = [=](const ptrdiff_t &v,
751  const ptrdiff_t &divisor) -> ptrdiff_t {
752  return ((v + divisor - 1) / divisor) * divisor;
753  };
754  typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
755  x;
756  typename detail::remove_pointer<OutPtr>::type carry = init;
757  for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
758  ptrdiff_t i = chunk + offset;
759  if (i < N) {
760  x = first[i];
761  }
763  exclusive_scan_over_group(g, x, carry, binary_op);
764  if (i < N) {
765  result[i] = out;
766  }
767  carry = group_broadcast(g, binary_op(out, x), stride - 1);
768  }
769  return result + N;
770 #else
771  (void)g;
772  (void)last;
773  (void)result;
774  (void)init;
775  throw runtime_error("Group algorithms are not supported on host device.",
776  PI_ERROR_INVALID_DEVICE);
777 #endif
778 }
779 
780 template <typename Group, typename InPtr, typename OutPtr,
781  class BinaryOperation>
783  (is_group_v<std::decay_t<Group>> && detail::is_pointer<InPtr>::value &&
784  detail::is_pointer<OutPtr>::value &&
786  typename detail::remove_pointer<InPtr>::type>::value &&
787  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
788  BinaryOperation>::value &&
789  detail::is_plus_if_complex<typename detail::remove_pointer<InPtr>::type,
790  BinaryOperation>::value),
791  OutPtr>
792 joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
793  BinaryOperation binary_op) {
794  // FIXME: Do not special-case for half precision
795  static_assert(
796  std::is_same<decltype(binary_op(*first, *first)),
797  typename detail::remove_pointer<OutPtr>::type>::value ||
798  (std::is_same<typename detail::remove_pointer<OutPtr>::type,
799  half>::value &&
800  std::is_same<decltype(binary_op(*first, *first)), float>::value),
801  "Result type of binary_op must match scan accumulation type.");
802  using T = typename detail::remove_pointer<InPtr>::type;
803  T init = detail::identity_for_ga_op<T, BinaryOperation>();
804  return joint_exclusive_scan(g, first, last, result, init, binary_op);
805 }
806 
807 // ---- inclusive_scan_over_group
808 // this function has two overloads, one with three arguments and one with four
809 // the three argument version is specialized thrice: vector, scalar, and
810 // complex
811 template <typename Group, typename T, class BinaryOperation>
812 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
813  detail::is_vector_arithmetic<T>::value &&
814  detail::is_native_op<T, BinaryOperation>::value),
815  T>
816 inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
817  // FIXME: Do not special-case for half precision
818  static_assert(
819  std::is_same<decltype(binary_op(x[0], x[0])),
820  typename T::element_type>::value ||
821  (std::is_same<T, half>::value &&
822  std::is_same<decltype(binary_op(x[0], x[0])), float>::value),
823  "Result type of binary_op must match scan accumulation type.");
824  T result;
825  for (int s = 0; s < x.size(); ++s) {
826  result[s] = inclusive_scan_over_group(g, x[s], binary_op);
827  }
828  return result;
829 }
830 
831 template <typename Group, typename T, class BinaryOperation>
832 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
833  detail::is_scalar_arithmetic<T>::value &&
834  detail::is_native_op<T, BinaryOperation>::value),
835  T>
836 inclusive_scan_over_group(Group, T x, BinaryOperation binary_op) {
837  // FIXME: Do not special-case for half precision
838  static_assert(std::is_same<decltype(binary_op(x, x)), T>::value ||
839  (std::is_same<T, half>::value &&
840  std::is_same<decltype(binary_op(x, x)), float>::value),
841  "Result type of binary_op must match scan accumulation type.");
842 #ifdef __SYCL_DEVICE_ONLY__
843  return sycl::detail::calc<T, __spv::GroupOperation::InclusiveScan,
844  sycl::detail::spirv::group_scope<Group>::value>(
845  typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
846 #else
847  throw runtime_error("Group algorithms are not supported on host device.",
848  PI_ERROR_INVALID_DEVICE);
849 #endif
850 }
851 
852 // complex specializaiton
853 template <typename Group, typename T, class BinaryOperation>
854 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
855  detail::is_complex<T>::value &&
856  detail::is_native_op<T, sycl::plus<T>>::value &&
857  detail::is_plus<T, BinaryOperation>::value),
858  T>
859 inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
860 #ifdef __SYCL_DEVICE_ONLY__
861  T result;
862  result.real(inclusive_scan_over_group(g, x.real(), sycl::plus<>()));
863  result.imag(inclusive_scan_over_group(g, x.imag(), sycl::plus<>()));
864  return result;
865 #else
866  (void)g;
867  (void)x;
868  (void)binary_op;
869  throw runtime_error("Group algorithms are not supported on host device.",
870  PI_ERROR_INVALID_DEVICE);
871 #endif
872 }
873 
874 // four argument version of inclusive_scan_over_group is specialized twice
875 // once for (scalar_arithmetic || complex) and once for vector_arithmetic
876 template <typename Group, typename V, class BinaryOperation, typename T>
878  (is_group_v<std::decay_t<Group>> &&
879  (detail::is_scalar_arithmetic<V>::value || detail::is_complex<V>::value) &&
880  (detail::is_scalar_arithmetic<T>::value || detail::is_complex<T>::value) &&
881  detail::is_native_op<V, BinaryOperation>::value &&
882  detail::is_native_op<T, BinaryOperation>::value &&
883  detail::is_plus_if_complex<T, BinaryOperation>::value &&
884  detail::is_plus_if_complex<V, BinaryOperation>::value),
885  T>
886 inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
887  // FIXME: Do not special-case for half precision
888  static_assert(std::is_same<decltype(binary_op(init, x)), T>::value ||
889  (std::is_same<T, half>::value &&
890  std::is_same<decltype(binary_op(init, x)), float>::value),
891  "Result type of binary_op must match scan accumulation type.");
892 #ifdef __SYCL_DEVICE_ONLY__
893  if (sycl::detail::get_local_linear_id(g) == 0) {
894  x = binary_op(init, x);
895  }
896  return inclusive_scan_over_group(g, x, binary_op);
897 #else
898  (void)g;
899  throw runtime_error("Group algorithms are not supported on host device.",
900  PI_ERROR_INVALID_DEVICE);
901 #endif
902 }
903 
904 template <typename Group, typename V, class BinaryOperation, typename T>
905 detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
906  detail::is_vector_arithmetic<V>::value &&
907  detail::is_vector_arithmetic<T>::value &&
908  detail::is_native_op<V, BinaryOperation>::value &&
909  detail::is_native_op<T, BinaryOperation>::value),
910  T>
911 inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
912  // FIXME: Do not special-case for half precision
913  static_assert(
914  std::is_same<decltype(binary_op(init[0], x[0])), T>::value ||
915  (std::is_same<T, half>::value &&
916  std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
917  "Result type of binary_op must match scan accumulation type.");
918  T result;
919  for (int s = 0; s < x.size(); ++s) {
920  result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]);
921  }
922  return result;
923 }
924 
925 // ---- joint_inclusive_scan
926 template <typename Group, typename InPtr, typename OutPtr,
927  class BinaryOperation, typename T>
929  (is_group_v<std::decay_t<Group>> && detail::is_pointer<InPtr>::value &&
930  detail::is_pointer<OutPtr>::value &&
932  typename detail::remove_pointer<InPtr>::type>::value &&
933  detail::is_arithmetic_or_complex<T>::value &&
934  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
935  BinaryOperation>::value &&
936  detail::is_native_op<T, BinaryOperation>::value &&
937  detail::is_plus_if_complex<typename detail::remove_pointer<InPtr>::type,
938  BinaryOperation>::value &&
939  detail::is_plus_if_complex<T, BinaryOperation>::value),
940  OutPtr>
941 joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
942  BinaryOperation binary_op, T init) {
943  // FIXME: Do not special-case for half precision
944  static_assert(
945  std::is_same<decltype(binary_op(init, *first)), T>::value ||
946  (std::is_same<T, half>::value &&
947  std::is_same<decltype(binary_op(init, *first)), float>::value),
948  "Result type of binary_op must match scan accumulation type.");
949 #ifdef __SYCL_DEVICE_ONLY__
950  ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
951  ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
952  ptrdiff_t N = last - first;
953  auto roundup = [=](const ptrdiff_t &v,
954  const ptrdiff_t &divisor) -> ptrdiff_t {
955  return ((v + divisor - 1) / divisor) * divisor;
956  };
957  typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
958  x;
959  typename detail::remove_pointer<OutPtr>::type carry = init;
960  for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
961  ptrdiff_t i = chunk + offset;
962  if (i < N) {
963  x = first[i];
964  }
966  inclusive_scan_over_group(g, x, binary_op, carry);
967  if (i < N) {
968  result[i] = out;
969  }
970  carry = group_broadcast(g, out, stride - 1);
971  }
972  return result + N;
973 #else
974  (void)g;
975  (void)last;
976  (void)result;
977  throw runtime_error("Group algorithms are not supported on host device.",
978  PI_ERROR_INVALID_DEVICE);
979 #endif
980 }
981 
982 template <typename Group, typename InPtr, typename OutPtr,
983  class BinaryOperation>
985  (is_group_v<std::decay_t<Group>> && detail::is_pointer<InPtr>::value &&
986  detail::is_pointer<OutPtr>::value &&
988  typename detail::remove_pointer<InPtr>::type>::value &&
989  detail::is_native_op<typename detail::remove_pointer<InPtr>::type,
990  BinaryOperation>::value &&
991  detail::is_plus_if_complex<typename detail::remove_pointer<InPtr>::type,
992  BinaryOperation>::value),
993  OutPtr>
994 joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
995  BinaryOperation binary_op) {
996  // FIXME: Do not special-case for half precision
997  static_assert(
998  std::is_same<decltype(binary_op(*first, *first)),
999  typename detail::remove_pointer<OutPtr>::type>::value ||
1000  (std::is_same<typename detail::remove_pointer<OutPtr>::type,
1001  half>::value &&
1002  std::is_same<decltype(binary_op(*first, *first)), float>::value),
1003  "Result type of binary_op must match scan accumulation type.");
1004 
1005  using T = typename detail::remove_pointer<InPtr>::type;
1006  T init = detail::identity_for_ga_op<T, BinaryOperation>();
1007  return joint_inclusive_scan(g, first, last, result, binary_op, init);
1008 }
1009 
1010 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1011 } // namespace sycl
A unique identifier of an item in an index space.
Definition: id.hpp:32
#define __SYCL_INLINE_VER_NAMESPACE(X)
std::integral_constant< bool,(is_complex< T >::value ? is_plus< T, BinaryOperation >::value :std::true_type::value)> is_plus_if_complex
std::integral_constant< bool, sycl::detail::is_complex< T >::value||sycl::detail::is_arithmetic< T >::value > is_arithmetic_or_complex
std::integral_constant< bool, std::is_same< BinaryOperation, sycl::plus< T > >::value||std::is_same< BinaryOperation, sycl::plus< void > >::value > is_plus
id< 3 > linear_id_to_id(range< 3 > r, size_t linear_id)
typename std::enable_if< B, T >::type enable_if_t
size_t get_local_linear_range(Group g)
constexpr detail::enable_if_t<!is_complex< T >::value, T > identity_for_ga_op()
Group::linear_id_type get_local_linear_id(Group g)
Function for_each(Group g, Ptr first, Ptr last, Function f)
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 >> &&(std::is_trivially_copyable< T >::value||detail::is_vec< T >::value)), T > group_broadcast(Group g, T x)
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<(std::is_same< std::decay_t< Group >, sub_group >::value &&(std::is_trivially_copyable< T >::value||detail::is_vec< T >::value)), T > select_from_group(Group, T x, typename Group::id_type local_id)
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<(std::is_same< std::decay_t< Group >, sub_group >::value &&(std::is_trivially_copyable< T >::value||detail::is_vec< T >::value)), T > shift_group_left(Group, T x, typename Group::linear_id_type delta=1)
detail::enable_if_t<(std::is_same< std::decay_t< Group >, sub_group >::value &&(std::is_trivially_copyable< T >::value||detail::is_vec< T >::value)), T > shift_group_right(Group, T x, typename Group::linear_id_type delta=1)
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<(std::is_same< std::decay_t< Group >, sub_group >::value &&(std::is_trivially_copyable< T >::value||detail::is_vec< T >::value)), T > permute_group_by_xor(Group, T x, typename Group::linear_id_type mask)
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)
sycl::detail::half_impl::half half
Definition: aliases.hpp:99
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
range_type get_local_range() const
Definition: sub_group.hpp:160