Skip to content

Clarify group algorithm constraints #730

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Pennycook wants to merge 4 commits into
base: main
Choose a base branch
from
Open

Clarify group algorithm constraints #730

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1aa0807
Further clarify group algorithm constraints
Pennycook Feb 18, 2025
1454481
Add sycl::half to group algorithm constraints
Pennycook Feb 18, 2025
ea2e1c9
Remove note about vec/marray for group functions
Pennycook Feb 18, 2025
cd2801c
Clarify group algorithms expect fundamental marray
Pennycook Mar 10, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
195 changes: 123 additions & 72 deletions adoc/chapters/programming_interface.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <<table.types.fundamental>>) 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
Expand Down Expand Up @@ -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 <<table.types.fundamental>>) and instances of the SYCL [code]#vec# and
[code]#marray# classes.

The <<group>> 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 <<sec:group-functions>>),
Expand Down Expand Up @@ -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<std::decay_t<Group>># is true and [code]#Ptr# is a
pointer.
. _Constraints:_ Available only if all of the following conditions are met:
- [code]#sycl::is_group_v<std::decay_t<Group>># is true; and
- [code]#Ptr# is a pointer or [code]#multi_ptr#.

+
--
_Preconditions:_ [code]#first# and [code]#last# must be the same for all
Expand Down Expand Up @@ -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<std::decay_t<Group>># is true and [code]#Ptr# is a
pointer.
. _Constraints:_ Available only if all of the following conditions are met:
- [code]#sycl::is_group_v<std::decay_t<Group>># is true; and
- [code]#Ptr# is a pointer or [code]#multi_ptr#.

+
--
_Preconditions:_ [code]#first# and [code]#last# must be the same for all
Expand Down Expand Up @@ -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<std::decay_t<Group>># is true and [code]#Ptr# is a
pointer.
. _Constraints:_ Available only if all of the following conditions are met:
- [code]#sycl::is_group_v<std::decay_t<Group>># is true; and
- [code]#Ptr# is a pointer or [code]#multi_ptr#.

+
--
_Preconditions:_ [code]#first# and [code]#last# must be the same for all
Expand Down Expand Up @@ -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<std::decay_t<Group>,
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<std::decay_t<Group>, 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.
Expand All @@ -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<std::decay_t<Group>,
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<std::decay_t<Group>, 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.
Expand Down Expand Up @@ -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<std::decay_t<Group>,
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<std::decay_t<Group>, 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.
Expand Down Expand Up @@ -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<std::decay_t<Group>,
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<std::decay_t<Group>, 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
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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
Expand All @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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
Expand All @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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#.
Expand All @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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#.
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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#.
Expand All @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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#.
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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
Expand Down Expand Up @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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#.
Expand All @@ -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<std::decay_t<Group>># 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<std::decay_t<Group>># 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#.
Expand Down