From 1aa080789d71f97be4b76d50238cd9e618770767 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 18 Feb 2025 10:19:53 +0000 Subject: [PATCH 1/4] Further clarify group algorithm constraints This commit makes two clarifications: - When an algorithm accepts a pointer, it can be a multi_ptr; and - When an algorithm accepts a value, it can be a vec or marray. A list of types compatible with group algorithms is removed in favor of listing the types explicitly for each algorithm, since some algorithms have different constraints. --- adoc/chapters/programming_interface.adoc | 176 ++++++++++++++--------- 1 file changed, 108 insertions(+), 68 deletions(-) diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 67d9ce8e0..4e073150c 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -21712,10 +21712,6 @@ Some of the functions in the SYCL algorithms library introduce additional restrictions in order to maximize portability across different devices and to minimize the chances of encountering unexpected behavior. -All algorithms are supported for the fundamental scalar types supported by SYCL -(see <>) and instances of the SYCL [code]#vec# and -[code]#marray# classes. - The <> argument to a SYCL algorithm denotes that it should be performed collaboratively by the work-items in the specified group. All algorithms act as group functions (as defined in <>), @@ -21752,9 +21748,10 @@ Boolean conditions applied to data held directly by the work-items in a group. include::{header_dir}/algorithms/any_of.h[lines=4..-1] ---- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true and [code]#Ptr# is a - pointer. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; and + - [code]#Ptr# is a pointer or [code]#multi_ptr#. + + -- _Preconditions:_ [code]#first# and [code]#last# must be the same for all @@ -21812,9 +21809,10 @@ _Returns:_ true if [code]#pred# is true for any work-item in group [code]#g#. include::{header_dir}/algorithms/all_of.h[lines=4..-1] ---- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true and [code]#Ptr# is a - pointer. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; and + - [code]#Ptr# is a pointer or [code]#multi_ptr#. + + -- _Preconditions:_ [code]#first# and [code]#last# must be the same for all @@ -21872,9 +21870,10 @@ _Returns:_ true if [code]#pred# is true for all work-items in group [code]#g#. include::{header_dir}/algorithms/none_of.h[lines=4..-1] ---- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true and [code]#Ptr# is a - pointer. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; and + - [code]#Ptr# is a pointer or [code]#multi_ptr#. + + -- _Preconditions:_ [code]#first# and [code]#last# must be the same for all @@ -21943,8 +21942,10 @@ shifting values a fixed number of work-items to the left or right. include::{header_dir}/algorithms/shift.h[lines=4..-1] ---- - . _Constraints:_ Available only if [code]#std::is_same_v, - sub_group># is true and [code]#T# is a trivially copyable type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_same_v, sub_group># is true; and + - [code]#T# is a trivially copyable type. + + -- _Preconditions:_ [code]#delta# must be the same for all work-items in the group. @@ -21963,8 +21964,10 @@ _Returns:_ the value of [code]#x# from the work-item whose group local id size, but the value returned in this case is unspecified. -- - . _Constraints:_ Available only if [code]#std::is_same_v, - sub_group># is true and [code]#T# is a trivially copyable type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_same_v, sub_group># is true; and + - [code]#T# is a trivially copyable type. + + -- _Preconditions:_ [code]#delta# must be the same for all work-items in the group. @@ -21997,8 +22000,10 @@ id and some fixed mask. include::{header_dir}/algorithms/permute.h[lines=4..-1] ---- - . _Constraints:_ Available only if [code]#std::is_same_v, - sub_group># is true and [code]#T# is a trivially copyable type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_same_v, sub_group># is true; and + - [code]#T# is a trivially copyable type. + + -- _Preconditions:_ [code]#mask# must be the same for all work-items in the group. @@ -22031,8 +22036,10 @@ by any other work-item in group [code]#g#. include::{header_dir}/algorithms/select.h[lines=4..-1] ---- - . _Constraints:_ Available only if [code]#std::is_same_v, - sub_group># is true and [code]#T# is a trivially copyable type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_same_v, sub_group># is true; and + - [code]#T# is a trivially copyable type. + + -- _Effects:_ Blocks until all work-items in group [code]#g# have reached this @@ -22074,10 +22081,12 @@ used for forward compatibility with future SYCL versions. include::{header_dir}/algorithms/reduce.h[lines=4..-1] ---- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#Ptr# is a - pointer to a fundamental type, and [code]#BinaryOperation# is a SYCL - function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#Ptr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(*first, *first)# must return a value of type @@ -22101,10 +22110,13 @@ iterators in the range [code]#[first, last)# using the operator sum defined in standard {cpp}. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#Ptr# is a - pointer to a fundamental type, [code]#T# is a fundamental type, and - [code]#BinaryOperation# is a SYCL function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#Ptr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(init, *first)# must return a value of type @@ -22128,9 +22140,11 @@ using the operator [code]#binary_op#, where the values are combined according to the generalized sum defined in standard {cpp}. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#T# is a - fundamental type and [code]#BinaryOperation# is a SYCL function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(x, x)# must return a value of type [code]#T#. @@ -22151,10 +22165,12 @@ work-item in group [code]#g# using the operator [code]#binary_op#, where the values are combined according to the generalized sum defined in standard {cpp}. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#V# and - [code]#T# are fundamental types, and [code]#BinaryOperation# is a SYCL - function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#V# is a fundamental type, [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(init, x)# must return a value of type [code]#T#. @@ -22208,10 +22224,14 @@ forward compatibility with future SYCL versions. include::{header_dir}/algorithms/exclusive_scan.h[lines=4..-1] ---- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#InPtr# and - [code]#OutPtr# are pointers to fundamental types, and - [code]#BinaryOperation# is a SYCL function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(*first, *first)# must return a value of type @@ -22244,10 +22264,15 @@ same synchronization point is unblocked. _Returns:_ A pointer to the end of the output range. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#InPtr# and - [code]#OutPtr# are pointers to fundamental types, [code]#T# is a fundamental - type, and [code]#BinaryOperation# is a SYCL function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(init, *first)# must return a value of type @@ -22281,10 +22306,11 @@ same synchronization point is unblocked. _Returns:_ A pointer to the end of the output range. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#T# is a - fundamental type, and [code]#BinaryOperation# is a SYCL function object - type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(x, x)# must return a value of type [code]#T#. @@ -22310,10 +22336,12 @@ For multi-dimensional groups, the order of work-items in group [code]#g# is determined by their linear id. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#V# and - [code]#T# are fundamental types, and [code]#BinaryOperation# is a SYCL - function object type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#V# is a fundamental type, [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(init, x)# must return a value of type [code]#T#. @@ -22343,10 +22371,14 @@ determined by their linear id. include::{header_dir}/algorithms/inclusive_scan.h[lines=4..-1] ---- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#InPtr# and - [code]#OutPtr# are pointers to fundamental types, and - [code]#BinaryOperation# is a SYCL function object type. +. _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(*first, *first)# must return a value of type @@ -22378,10 +22410,15 @@ same synchronization point is unblocked. _Returns:_ A pointer to the end of the output range. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#InPtr# and - [code]#OutPtr# are pointers to fundamental types, [code]#BinaryOperation# is - a SYCL function object type, and [code]#T# is a fundamental type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, + [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#BinaryOperation# is a SYCL function object type. + + -- _Mandates:_ [code]#binary_op(init, *first)# must return a value of type @@ -22415,10 +22452,11 @@ same synchronization point is unblocked. _Returns:_ A pointer to the end of the output range. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#T# is a - fundamental type, and [code]#BinaryOperation# is a SYCL function object - type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#BinaryOperation# is a SYCL function object type; and + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#. + + -- _Mandates:_ [code]#binary_op(x, x)# must return a value of type [code]#T#. @@ -22442,10 +22480,12 @@ For multi-dimensional groups, the order of work-items in group [code]#g# is determined by their linear id. -- - . _Constraints:_ Available only if - [code]#sycl::is_group_v># is true, [code]#V# is a - fundamental type, [code]#BinaryOperation# is a SYCL function object type, - and [code]#T# is a fundamental type. + . _Constraints:_ Available only if all of the following conditions are met: + - [code]#sycl::is_group_v># is true; + - [code]#V# is a fundamental type, [code]#marray# or [code]#vec#; + - [code]#BinaryOperation# is a SYCL function object type; and + - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#. + + -- _Mandates:_ [code]#binary_op(init, x)# must return a value of type [code]#T#. From 1454481d1c6fa11044eb89ba74237ef0875ca50a Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 18 Feb 2025 10:30:40 +0000 Subject: [PATCH 2/4] Add sycl::half to group algorithm constraints --- adoc/chapters/programming_interface.adoc | 56 ++++++++++++++---------- 1 file changed, 34 insertions(+), 22 deletions(-) diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 4e073150c..3739c3628 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -22084,7 +22084,7 @@ include::{header_dir}/algorithms/reduce.h[lines=4..-1] . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#Ptr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22113,8 +22113,9 @@ sum defined in standard {cpp}. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#Ptr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22142,7 +22143,8 @@ the generalized sum defined in standard {cpp}. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22167,8 +22169,10 @@ values are combined according to the generalized sum defined in standard {cpp}. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#V# is a fundamental type, [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22227,9 +22231,9 @@ include::{header_dir}/algorithms/exclusive_scan.h[lines=4..-1] . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22267,10 +22271,11 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22308,7 +22313,8 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22338,8 +22344,10 @@ determined by their linear id. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#V# is a fundamental type, [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22374,9 +22382,9 @@ include::{header_dir}/algorithms/inclusive_scan.h[lines=4..-1] . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22413,10 +22421,11 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22455,7 +22464,8 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#BinaryOperation# is a SYCL function object type; and - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#. + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#. + -- @@ -22482,9 +22492,11 @@ determined by their linear id. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#V# is a fundamental type, [code]#marray# or [code]#vec#; + - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#; - [code]#BinaryOperation# is a SYCL function object type; and - - [code]#T# is a fundamental type, [code]#marray# or [code]#vec#. + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or + [code]#vec#. + -- From ea2e1c9b9862860cffc30dc030b8e4e0e56c3d3e Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 18 Feb 2025 14:27:52 +0000 Subject: [PATCH 3/4] Remove note about vec/marray for group functions This note had been duplicated for group functions and group algorithms. The note about group algorithms was removed in a previous commit. --- adoc/chapters/programming_interface.adoc | 4 ---- 1 file changed, 4 deletions(-) diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 3739c3628..ed7d1dcc5 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -21571,10 +21571,6 @@ Any such restrictions on the arguments passed to a function are defined within the descriptions of those functions. Violating these restrictions results in undefined behavior. -All group functions are supported for the fundamental scalar types supported by -SYCL (see <>) and instances of the SYCL [code]#vec# and -[code]#marray# classes. - Using a group function inside of a kernel may introduce additional limits on the resources available to user code inside the same kernel. The behavior of these limits is implementation-defined, but must be reflected by From cd2801cd02b3cd8d7c8261099e46eeecd3223eb7 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Mon, 10 Mar 2025 10:01:57 +0000 Subject: [PATCH 4/4] Clarify group algorithms expect fundamental marray --- adoc/chapters/programming_interface.adoc | 71 ++++++++++++------------ 1 file changed, 37 insertions(+), 34 deletions(-) diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index ed7d1dcc5..a0416b30c 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -22080,7 +22080,8 @@ include::{header_dir}/algorithms/reduce.h[lines=4..-1] . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#Ptr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; + and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22109,9 +22110,9 @@ sum defined in standard {cpp}. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#Ptr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22139,8 +22140,8 @@ the generalized sum defined in standard {cpp}. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22165,10 +22166,10 @@ values are combined according to the generalized sum defined in standard {cpp}. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22227,9 +22228,10 @@ include::{header_dir}/algorithms/exclusive_scan.h[lines=4..-1] . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; + and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22267,11 +22269,11 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22309,8 +22311,8 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22340,10 +22342,10 @@ determined by their linear id. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22378,9 +22380,10 @@ include::{header_dir}/algorithms/inclusive_scan.h[lines=4..-1] . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; and + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; + and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22417,11 +22420,11 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#InPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; - [code]#OutPtr# is a pointer or [code]#multi_ptr# to: a fundamental type, - [code]#sycl::half#, [code]#marray# or [code]#vec#; - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; and + [code]#sycl::half#, [code]#marray# of a fundamental type, or [code]#vec#; + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; and - [code]#BinaryOperation# is a SYCL function object type. + @@ -22460,8 +22463,8 @@ _Returns:_ A pointer to the end of the output range. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - [code]#BinaryOperation# is a SYCL function object type; and - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#. + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#. + -- @@ -22488,11 +22491,11 @@ determined by their linear id. . _Constraints:_ Available only if all of the following conditions are met: - [code]#sycl::is_group_v># is true; - - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#; + - [code]#V# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#; - [code]#BinaryOperation# is a SYCL function object type; and - - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# or - [code]#vec#. + - [code]#T# is a fundamental type, [code]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#. + --