From 49e4962cb65ee9b00746a5f44a03ebf6b93713ad Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 10 Sep 2020 14:54:50 -0400 Subject: [PATCH 01/12] [SYCL][Doc] Replace is_native_function_object is_native_function_object proved too difficult to implement and use in a practical way. Replacing it with constraints makes the behavior of group algorithms more predictable and enables users to detect which type combinations are supported at compile-time. Signed-off-by: John Pennycook --- .../SYCL_INTEL_group_algorithms.asciidoc | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc index 31392cd740a82..dff55b956e16b 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc @@ -118,7 +118,7 @@ The following function objects alias objects in the ++ header from t - +cl::sycl::intel::bit_xor+ - +cl::sycl::intel::logical_and+ - +cl::sycl::intel::logical_or+ - +/ New function objects without {cpp} standard library equivalents are defined in the table below: |=== @@ -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. +All restrictions on acceptable group types, input types and function objects must be implemented as constraints. + ==== 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* |======================================== //************************************************************************ From 17582b3a4b06c891825214eeef7cc4fb2e2570b1 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 10 Sep 2020 16:29:07 -0400 Subject: [PATCH 02/12] [SYCL] Add constraints to all_of/any_of/none_of Signed-off-by: John Pennycook --- .../CL/sycl/ONEAPI/group_algorithm.hpp | 75 +++++++++---------- 1 file changed, 36 insertions(+), 39 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index d93b068c09378..792bf9d23c0b2 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,19 @@ 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<(detail::is_generic_group::value && + detail::is_arithmetic::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 && + detail::is_arithmetic::type>::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 +232,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 +245,20 @@ 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<(detail::is_generic_group::value && + detail::is_arithmetic::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 && + detail::is_arithmetic::type>::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 +274,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 +287,20 @@ 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<(detail::is_generic_group::value && + detail::is_arithmetic::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 && + detail::is_arithmetic::type>::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; From 658bca636d4b2eda91d1beddba88893b76e6c605 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 10 Sep 2020 16:48:42 -0400 Subject: [PATCH 03/12] [SYCL] Add constraints to broadcast Signed-off-by: John Pennycook --- .../CL/sycl/ONEAPI/group_algorithm.hpp | 51 +++++++++---------- 1 file changed, 25 insertions(+), 26 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index 792bf9d23c0b2..4fd874898b6f0 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -313,11 +313,11 @@ none_of(Group g, Ptr first, Ptr last, Predicate pred) { } 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 @@ -329,11 +329,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) { @@ -350,11 +349,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, @@ -369,11 +368,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) { @@ -390,10 +388,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 @@ -405,10 +404,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) { From 2b4848c82b33bbee5dd992acc6d01ae0c88686da Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 10 Sep 2020 17:26:01 -0400 Subject: [PATCH 04/12] [SYCL] Add constraints to reduce Signed-off-by: John Pennycook --- .../CL/sycl/ONEAPI/group_algorithm.hpp | 103 ++++++++++-------- 1 file changed, 57 insertions(+), 46 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index 4fd874898b6f0..6b11a8589f10f 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -423,11 +423,11 @@ 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 || @@ -445,11 +445,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) { +detail::enable_if_t<(detail::is_sub_group::value && + ((!detail::is_arithmetic::value && + std::is_trivially_copyable::value) || + !detail::is_native_op::value)), + T> +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."); @@ -481,11 +485,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 || @@ -502,11 +508,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 && + ((!(detail::is_arithmetic::value && + detail::is_arithmetic::value) && + (std::is_trivially_copyable::value && + std::is_trivially_copyable::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)); @@ -544,26 +555,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; @@ -575,11 +583,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 || @@ -587,12 +599,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; From 3c80f0fb8a3dca2a4fc5fe57d93913abcab9656f Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 10 Sep 2020 17:33:31 -0400 Subject: [PATCH 05/12] [SYCL] Add constraints to exclusive_scan Signed-off-by: John Pennycook --- .../CL/sycl/ONEAPI/group_algorithm.hpp | 63 ++++++++++++------- 1 file changed, 40 insertions(+), 23 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index 6b11a8589f10f..a9ffff82de874 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -614,11 +614,11 @@ reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { } 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 && @@ -635,11 +635,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 && @@ -705,12 +709,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 || @@ -751,9 +761,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 Date: Thu, 10 Sep 2020 17:54:03 -0400 Subject: [PATCH 06/12] [SYCL] Add constraints to inclusive_scan Signed-off-by: John Pennycook --- .../CL/sycl/ONEAPI/group_algorithm.hpp | 63 ++++++++++++------- 1 file changed, 40 insertions(+), 23 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index a9ffff82de874..33eee85f90771 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -785,11 +785,11 @@ exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, } 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 && @@ -826,11 +826,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 && @@ -849,11 +851,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 || @@ -869,12 +873,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 || @@ -914,9 +924,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 Date: Thu, 10 Sep 2020 17:57:12 -0400 Subject: [PATCH 07/12] [SYCL] Add constraints to leader Signed-off-by: John Pennycook --- sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index 33eee85f90771..f7efba372f951 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -946,10 +946,9 @@ inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, 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); From 66a13c06e573053ee28a73c451b00bc253b6b0d4 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 11 Sep 2020 12:14:24 -0400 Subject: [PATCH 08/12] [SYCL] Simplify sub-group constraints Conjunction of positive requirements && disjunction of negative requirements is easier to read and understand. Signed-off-by: John Pennycook --- sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index f7efba372f951..f9fa7b46dafbe 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -466,14 +466,11 @@ reduce(Group g, T x, BinaryOperation binary_op) { template detail::enable_if_t<(detail::is_sub_group::value && - ((!detail::is_arithmetic::value && - std::is_trivially_copyable::value) || + std::is_trivially_copyable::value && + (!detail::is_arithmetic::value || !detail::is_native_op::value)), T> 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."); 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)); @@ -537,10 +534,10 @@ reduce(Group g, V x, T init, BinaryOperation binary_op) { template detail::enable_if_t<(detail::is_sub_group::value && - ((!(detail::is_arithmetic::value && - detail::is_arithmetic::value) && - (std::is_trivially_copyable::value && - std::is_trivially_copyable::value)) || + std::is_trivially_copyable::value && + std::is_trivially_copyable::value && + (!std::is_arithmetic::value || + !std::is_arithmetic::value || !detail::is_native_op::value)), T> reduce(Group g, V x, T init, BinaryOperation op) { From b072471eb904188bffa68ccbc53b3b616b087557 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 11 Sep 2020 12:18:06 -0400 Subject: [PATCH 09/12] [SYCL][Doc] Remove accidental / Signed-off-by: John Pennycook --- .../GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc index dff55b956e16b..8bab4357cc4f3 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc @@ -118,7 +118,7 @@ The following function objects alias objects in the ++ header from t - +cl::sycl::intel::bit_xor+ - +cl::sycl::intel::logical_and+ - +cl::sycl::intel::logical_or+ -/ + New function objects without {cpp} standard library equivalents are defined in the table below: |=== From da518f920477a010766a6d7530594408ee79a9e4 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 11 Sep 2020 13:06:06 -0400 Subject: [PATCH 10/12] [SYCL] Use detail::is_arithmetic instead of std:: Should fix regression in fp16 reduction. Signed-off-by: John Pennycook --- sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index f9fa7b46dafbe..9cb88e53f13e4 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -536,8 +536,8 @@ template detail::enable_if_t<(detail::is_sub_group::value && std::is_trivially_copyable::value && std::is_trivially_copyable::value && - (!std::is_arithmetic::value || - !std::is_arithmetic::value || + (!detail::is_arithmetic::value || + !detail::is_arithmetic::value || !detail::is_native_op::value)), T> reduce(Group g, V x, T init, BinaryOperation op) { From 282634f7426200d35b8447afb5def2fa6b346876 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 18 Sep 2020 10:43:40 -0400 Subject: [PATCH 11/12] [SYCL] Remove any/all/none arithmetic restriction pred(x) will transform any input type to Boolean. Signed-off-by: John Pennycook --- .../CL/sycl/ONEAPI/group_algorithm.hpp | 33 +++++++------------ 1 file changed, 12 insertions(+), 21 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp index 9cb88e53f13e4..3aaab50b236aa 100644 --- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp +++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp @@ -203,18 +203,15 @@ all_of(Group, bool pred) { } template -detail::enable_if_t<(detail::is_generic_group::value && - detail::is_arithmetic::value), - bool> +detail::enable_if_t::value, bool> all_of(Group g, T x, Predicate pred) { return all_of(g, pred(x)); } template -detail::enable_if_t< - (detail::is_generic_group::value && detail::is_pointer::value && - detail::is_arithmetic::type>::value), - bool> +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; @@ -245,18 +242,15 @@ any_of(Group, bool pred) { } template -detail::enable_if_t<(detail::is_generic_group::value && - detail::is_arithmetic::value), - bool> +detail::enable_if_t::value, bool> any_of(Group g, T x, Predicate pred) { return any_of(g, pred(x)); } template -detail::enable_if_t< - (detail::is_generic_group::value && detail::is_pointer::value && - detail::is_arithmetic::type>::value), - bool> +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__ bool partial = false; @@ -287,18 +281,15 @@ none_of(Group, bool pred) { } template -detail::enable_if_t<(detail::is_generic_group::value && - detail::is_arithmetic::value), - bool> +detail::enable_if_t::value, bool> none_of(Group g, T x, Predicate pred) { return none_of(g, pred(x)); } template -detail::enable_if_t< - (detail::is_generic_group::value && detail::is_pointer::value && - detail::is_arithmetic::type>::value), - bool> +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__ return !any_of(g, first, last, pred); From c3165adb309b93853a9d2d053a762077aaf8f025 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 23 Sep 2020 13:31:55 -0700 Subject: [PATCH 12/12] [SYCL][Doc] Tweaked description of constraints Signed-off-by: John Pennycook --- .../GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc index 8bab4357cc4f3..6c6f65e9388d4 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc @@ -143,7 +143,7 @@ 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. -All restrictions on acceptable group types, input types and function objects must be implemented as constraints. +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