diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 67d9ce8e0..a0416b30c 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 @@ -21712,10 +21708,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 +21744,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 +21805,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 +21866,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 +21938,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 +21960,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 +21996,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 +22032,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 +22077,13 @@ 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]#sycl::half#, [code]#marray# of a fundamental type, 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 +22107,14 @@ 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]#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. + + -- _Mandates:_ [code]#binary_op(init, *first)# must return a value of type @@ -22128,9 +22138,12 @@ 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]#sycl::half#, [code]#marray# of a + fundamental type, 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 +22164,14 @@ 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]#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. + + -- _Mandates:_ [code]#binary_op(init, x)# must return a value of type [code]#T#. @@ -22208,10 +22225,15 @@ 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]#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# of a fundamental type, 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 +22266,16 @@ 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]#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# 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. + + -- _Mandates:_ [code]#binary_op(init, *first)# must return a value of type @@ -22281,10 +22309,12 @@ 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]#sycl::half#, [code]#marray# of a + fundamental type, 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 +22340,14 @@ 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]#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. + + -- _Mandates:_ [code]#binary_op(init, x)# must return a value of type [code]#T#. @@ -22343,10 +22377,15 @@ 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]#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# of a fundamental type, 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 +22417,16 @@ 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]#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# 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. + + -- _Mandates:_ [code]#binary_op(init, *first)# must return a value of type @@ -22415,10 +22460,12 @@ 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]#sycl::half#, [code]#marray# of a + fundamental type, or [code]#vec#. + + -- _Mandates:_ [code]#binary_op(x, x)# must return a value of type [code]#T#. @@ -22442,10 +22489,14 @@ 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]#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# of a + fundamental type, or [code]#vec#. + + -- _Mandates:_ [code]#binary_op(init, x)# must return a value of type [code]#T#.