diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc index 31392cd740a82..6c6f65e9388d4 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc @@ -131,8 +131,6 @@ New function objects without {cpp} standard library equivalents are defined in t |+T operator(const T&, const T&) const+ applies +std::greater+ to its arguments, in the same order, then returns the greater argument unchanged. |=== -Function objects supported by the group algorithms library can be identified using the +cl::sycl::intel::is_native_function_object+ and +cl::sycl::intel::is_native_function_object_v+ traits classes. - === Functions The group algorithms library is based on the algorithms library described in Section 28 of the {cpp}17 standard. The syntax and restrictions are aligned, with two notable differences: the first argument to each function is a group of work-items, in place of an execution policy; and pointers are accepted in place of iterators in order to guarantee that address space information is visible to the compiler. @@ -145,6 +143,8 @@ Using functions from the group algorithms library inside of a kernel may introdu It is undefined behavior for any of these functions to be invoked within a +parallel_for_work_group+ or +parallel_for_work_item+ context, but this restriction may be lifted in a future version of the proposal. +A number of the restrictions regarding the types of parameters that are acceptable for each algorithm must implemented as constraints: group arguments must be of a supported group class type; binary operations must be one of the group algorithms function objects; pointer arguments must be pointers to fundamental data types; and value arguments must be scalar fundamental data types (or vectors of those types). + ==== Vote |=== @@ -262,6 +262,11 @@ None. //*RESOLUTION*: Not resolved. //-- +. How should `is_native_function_object` work? Does it represent what is minimally required by the specification, or what the implementation really supports? +-- +*RESOLUTION*: The `is_native_function_object` trait has been removed. It proved too difficult to implement something that returned sensible values for transparent function objects (e.g. `std::plus`) that did not also require checking additional traits for each individual group algorithm. Requiring the user to implement their own checks based on type requirements outlined in the specification would make it significantly harder for implementers to extend the algorithms to types and function objects beyond what is specified. Using constrained forms of the algorithms instead allows a user to determine whether an implementation of a particular algorithm exists using the C++ detection idiom. +-- + == Revision History [cols="5,15,15,70"] @@ -270,6 +275,7 @@ None. |======================================== |Rev|Date|Author|Changes |1|2020-01-30|John Pennycook|*Initial public working draft* +|2|2020-09-10|John Pennycook|*Remove is_native_function_object and clarify which requirements are constraints* |======================================== //************************************************************************ diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index d93b068c09378..3aaab50b236aa 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -190,10 +190,9 @@ using EnableIfIsNonNativeOp = cl::sycl::detail::enable_if_t< !cl::sycl::detail::is_native_op::value, T>; -template bool all_of(Group, bool pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +template +detail::enable_if_t::value, bool> +all_of(Group, bool pred) { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::GroupAll(pred); #else @@ -204,19 +203,16 @@ template bool all_of(Group, bool pred) { } template -bool all_of(Group g, T x, Predicate pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t::value, bool> +all_of(Group g, T x, Predicate pred) { return all_of(g, pred(x)); } template -EnableIfIsPointer all_of(Group g, Ptr first, Ptr last, - Predicate pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_pointer::value), + bool> +all_of(Group g, Ptr first, Ptr last, Predicate pred) { #ifdef __SYCL_DEVICE_ONLY__ bool partial = true; sycl::detail::for_each( @@ -233,10 +229,9 @@ EnableIfIsPointer all_of(Group g, Ptr first, Ptr last, #endif } -template bool any_of(Group, bool pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +template +detail::enable_if_t::value, bool> +any_of(Group, bool pred) { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::GroupAny(pred); #else @@ -247,20 +242,17 @@ template bool any_of(Group, bool pred) { } template -bool any_of(Group g, T x, Predicate pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t::value, bool> +any_of(Group g, T x, Predicate pred) { return any_of(g, pred(x)); } template -EnableIfIsPointer any_of(Group g, Ptr first, Ptr last, - Predicate pred) { +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_pointer::value), + bool> +any_of(Group g, Ptr first, Ptr last, Predicate pred) { #ifdef __SYCL_DEVICE_ONLY__ - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); bool partial = false; sycl::detail::for_each( g, first, last, @@ -276,10 +268,9 @@ EnableIfIsPointer any_of(Group g, Ptr first, Ptr last, #endif } -template bool none_of(Group, bool pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +template +detail::enable_if_t::value, bool> +none_of(Group, bool pred) { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::GroupAll(!pred); #else @@ -290,20 +281,17 @@ template bool none_of(Group, bool pred) { } template -bool none_of(Group g, T x, Predicate pred) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t::value, bool> +none_of(Group g, T x, Predicate pred) { return none_of(g, pred(x)); } template -EnableIfIsPointer none_of(Group g, Ptr first, Ptr last, - Predicate pred) { +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_pointer::value), + bool> +none_of(Group g, Ptr first, Ptr last, Predicate pred) { #ifdef __SYCL_DEVICE_ONLY__ - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); return !any_of(g, first, last, pred); #else (void)g; @@ -316,11 +304,11 @@ EnableIfIsPointer none_of(Group g, Ptr first, Ptr last, } template -EnableIfIsTriviallyCopyable broadcast(Group, T x, - typename Group::id_type local_id) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_generic_group::value && + std::is_trivially_copyable::value && + !detail::is_vector_arithmetic::value), + T> +broadcast(Group, T x, typename Group::id_type local_id) { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::GroupBroadcast(x, local_id); #else @@ -332,11 +320,10 @@ EnableIfIsTriviallyCopyable broadcast(Group, T x, } template -EnableIfIsVectorArithmetic broadcast(Group g, T x, - typename Group::id_type local_id) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value), + T> +broadcast(Group g, T x, typename Group::id_type local_id) { #ifdef __SYCL_DEVICE_ONLY__ T result; for (int s = 0; s < x.get_size(); ++s) { @@ -353,11 +340,11 @@ EnableIfIsVectorArithmetic broadcast(Group g, T x, } template -EnableIfIsTriviallyCopyable +detail::enable_if_t<(detail::is_generic_group::value && + std::is_trivially_copyable::value && + !detail::is_vector_arithmetic::value), + T> broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); #ifdef __SYCL_DEVICE_ONLY__ return broadcast( g, x, @@ -372,11 +359,10 @@ broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) { } template -EnableIfIsVectorArithmetic +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value), + T> broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); #ifdef __SYCL_DEVICE_ONLY__ T result; for (int s = 0; s < x.get_size(); ++s) { @@ -393,10 +379,11 @@ broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) { } template -EnableIfIsTriviallyCopyable broadcast(Group g, T x) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_generic_group::value && + std::is_trivially_copyable::value && + !detail::is_vector_arithmetic::value), + T> +broadcast(Group g, T x) { #ifdef __SYCL_DEVICE_ONLY__ return broadcast(g, x, 0); #else @@ -408,10 +395,10 @@ EnableIfIsTriviallyCopyable broadcast(Group g, T x) { } template -EnableIfIsVectorArithmetic broadcast(Group g, T x) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value), + T> +broadcast(Group g, T x) { #ifdef __SYCL_DEVICE_ONLY__ T result; for (int s = 0; s < x.get_size(); ++s) { @@ -427,11 +414,11 @@ EnableIfIsVectorArithmetic broadcast(Group g, T x) { } template -EnableIfIsScalarArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_scalar_arithmetic::value && + detail::is_native_op::value), + T> reduce(Group, T x, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same::value || @@ -449,11 +436,11 @@ reduce(Group, T x, BinaryOperation binary_op) { } template -EnableIfIsVectorArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value && + detail::is_native_op::value), + T> reduce(Group g, T x, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same -EnableIfIsNonNativeOp reduce(Group g, T x, - BinaryOperation op) { - static_assert(sycl::detail::is_sub_group::value, - "reduce algorithm with user-defined types and operators" - "only supports ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_sub_group::value && + std::is_trivially_copyable::value && + (!detail::is_arithmetic::value || + !detail::is_native_op::value)), + T> +reduce(Group g, T x, BinaryOperation op) { T result = x; for (int mask = 1; mask < g.get_max_local_range()[0]; mask *= 2) { T tmp = g.shuffle_xor(result, id<1>(mask)); @@ -485,11 +473,13 @@ EnableIfIsNonNativeOp reduce(Group g, T x, } template -EnableIfIsScalarArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_scalar_arithmetic::value && + detail::is_scalar_arithmetic::value && + detail::is_native_op::value && + detail::is_native_op::value), + T> reduce(Group g, V x, T init, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same::value || @@ -506,11 +496,13 @@ reduce(Group g, V x, T init, BinaryOperation binary_op) { } template -EnableIfIsVectorArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value && + detail::is_vector_arithmetic::value && + detail::is_native_op::value && + detail::is_native_op::value), + T> reduce(Group g, V x, T init, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same -EnableIfIsNonNativeOp reduce(Group g, V x, T init, - BinaryOperation op) { - static_assert(sycl::detail::is_sub_group::value, - "reduce algorithm with user-defined types and operators" - "only supports ONEAPI::sub_group class."); +detail::enable_if_t<(detail::is_sub_group::value && + std::is_trivially_copyable::value && + std::is_trivially_copyable::value && + (!detail::is_arithmetic::value || + !detail::is_arithmetic::value || + !detail::is_native_op::value)), + T> +reduce(Group g, V x, T init, BinaryOperation op) { T result = x; for (int mask = 1; mask < g.get_max_local_range()[0]; mask *= 2) { T tmp = g.shuffle_xor(result, id<1>(mask)); @@ -548,26 +543,23 @@ EnableIfIsNonNativeOp reduce(Group g, V x, T init, } template -EnableIfIsPointer +detail::enable_if_t< + (detail::is_generic_group::value && detail::is_pointer::value && + detail::is_arithmetic::type>::value), + typename detail::remove_pointer::type> reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); + using T = typename detail::remove_pointer::type; // FIXME: Do not special-case for half precision static_assert( - std::is_same::value || - (std::is_same::value && + std::is_same::value || + (std::is_same::value && std::is_same::value), "Result type of binary_op must match reduction accumulation type."); #ifdef __SYCL_DEVICE_ONLY__ typename Ptr::element_type partial = - sycl::detail::identity::value; + sycl::detail::identity::value; sycl::detail::for_each(g, first, last, - [&](const typename Ptr::element_type &x) { - partial = binary_op(partial, x); - }); + [&](const T &x) { partial = binary_op(partial, x); }); return reduce(g, partial, binary_op); #else (void)g; @@ -579,11 +571,15 @@ reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) { } template -EnableIfIsPointer reduce(Group g, Ptr first, Ptr last, T init, - BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +detail::enable_if_t< + (detail::is_generic_group::value && detail::is_pointer::value && + detail::is_arithmetic::type>::value && + detail::is_arithmetic::value && + detail::is_native_op::type, + BinaryOperation>::value && + detail::is_native_op::value), + T> +reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { // FIXME: Do not special-case for half precision static_assert( std::is_same::value || @@ -591,12 +587,11 @@ EnableIfIsPointer reduce(Group g, Ptr first, Ptr last, T init, std::is_same::value), "Result type of binary_op must match reduction accumulation type."); #ifdef __SYCL_DEVICE_ONLY__ - T partial = sycl::detail::identity::value; - sycl::detail::for_each(g, first, last, - [&](const typename Ptr::element_type &x) { - partial = binary_op(partial, x); - }); + T partial = sycl::detail::identity::value; + sycl::detail::for_each( + g, first, last, [&](const typename detail::remove_pointer::type &x) { + partial = binary_op(partial, x); + }); return reduce(g, partial, init, binary_op); #else (void)g; @@ -607,11 +602,11 @@ EnableIfIsPointer reduce(Group g, Ptr first, Ptr last, T init, } template -EnableIfIsScalarArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_scalar_arithmetic::value && + detail::is_native_op::value), + T> exclusive_scan(Group, T x, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert(std::is_same::value || (std::is_same::value && @@ -628,11 +623,11 @@ exclusive_scan(Group, T x, BinaryOperation binary_op) { } template -EnableIfIsVectorArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value && + detail::is_native_op::value), + T> exclusive_scan(Group g, T x, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same -EnableIfIsVectorArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value && + detail::is_vector_arithmetic::value && + detail::is_native_op::value && + detail::is_native_op::value), + T> exclusive_scan(Group g, V x, T init, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same -EnableIfIsScalarArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_scalar_arithmetic::value && + detail::is_scalar_arithmetic::value && + detail::is_native_op::value && + detail::is_native_op::value), + T> exclusive_scan(Group g, V x, T init, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert(std::is_same::value || (std::is_same::value && @@ -698,12 +697,18 @@ exclusive_scan(Group g, V x, T init, BinaryOperation binary_op) { template -EnableIfIsPointer +detail::enable_if_t< + (detail::is_generic_group::value && + detail::is_pointer::value && detail::is_pointer::value && + detail::is_arithmetic< + typename detail::remove_pointer::type>::value && + detail::is_arithmetic::value && + detail::is_native_op::type, + BinaryOperation>::value && + detail::is_native_op::value), + OutPtr> exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same::value || @@ -744,9 +749,16 @@ exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, template -EnableIfIsPointer exclusive_scan(Group g, InPtr first, - InPtr last, OutPtr result, - BinaryOperation binary_op) { +detail::enable_if_t< + (detail::is_generic_group::value && + detail::is_pointer::value && detail::is_pointer::value && + detail::is_arithmetic< + typename detail::remove_pointer::type>::value && + detail::is_native_op::type, + BinaryOperation>::value), + OutPtr> +exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + BinaryOperation binary_op) { // FIXME: Do not special-case for half precision static_assert( std::is_same exclusive_scan(Group g, InPtr first, } template -EnableIfIsVectorArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value && + detail::is_native_op::value), + T> inclusive_scan(Group g, T x, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same -EnableIfIsScalarArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_scalar_arithmetic::value && + detail::is_native_op::value), + T> inclusive_scan(Group, T x, BinaryOperation binary_op) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert(std::is_same::value || (std::is_same::value && @@ -802,11 +814,13 @@ inclusive_scan(Group, T x, BinaryOperation binary_op) { } template -EnableIfIsScalarArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_scalar_arithmetic::value && + detail::is_scalar_arithmetic::value && + detail::is_native_op::value && + detail::is_native_op::value), + T> inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert(std::is_same::value || (std::is_same::value && @@ -825,11 +839,13 @@ inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) { } template -EnableIfIsVectorArithmeticNativeOp +detail::enable_if_t<(detail::is_generic_group::value && + detail::is_vector_arithmetic::value && + detail::is_vector_arithmetic::value && + detail::is_native_op::value && + detail::is_native_op::value), + T> inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same::value || @@ -845,12 +861,18 @@ inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) { template -EnableIfIsPointer +detail::enable_if_t< + (detail::is_generic_group::value && + detail::is_pointer::value && detail::is_pointer::value && + detail::is_arithmetic< + typename detail::remove_pointer::type>::value && + detail::is_arithmetic::value && + detail::is_native_op::type, + BinaryOperation>::value && + detail::is_native_op::value), + OutPtr> inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); // FIXME: Do not special-case for half precision static_assert( std::is_same::value || @@ -890,9 +912,16 @@ inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, template -EnableIfIsPointer inclusive_scan(Group g, InPtr first, - InPtr last, OutPtr result, - BinaryOperation binary_op) { +detail::enable_if_t< + (detail::is_generic_group::value && + detail::is_pointer::value && detail::is_pointer::value && + detail::is_arithmetic< + typename detail::remove_pointer::type>::value && + detail::is_native_op::type, + BinaryOperation>::value), + OutPtr> +inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + BinaryOperation binary_op) { // FIXME: Do not special-case for half precision static_assert( std::is_same inclusive_scan(Group g, InPtr first, BinaryOperation>::value); } -template bool leader(Group g) { - static_assert(sycl::detail::is_generic_group::value, - "Group algorithms only support the sycl::group and " - "ONEAPI::sub_group class."); +template +detail::enable_if_t::value, bool> +leader(Group g) { #ifdef __SYCL_DEVICE_ONLY__ typename Group::linear_id_type linear_id = sycl::detail::get_local_linear_id(g);