-
Notifications
You must be signed in to change notification settings - Fork 789
[SYCL][Doc] Add draft of sycl_ext_oneapi_joint_for #14886
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.
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
Merged
Merged
Changes from all commits
Commits
Show all changes
5 commits
Select commit
Hold shift + click to select a range
fa5564f
[SYCL][Doc] Add draft of sycl_ext_oneapi_joint_for
Pennycook 3669d47
Bump SYCL 2020 revision
Pennycook 301db76
Move extension to experimental namespace
Pennycook 34a6d88
Switch return type to void
Pennycook 1c3c61f
Add usage examples
Pennycook File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
261 changes: 261 additions & 0 deletions
261
sycl/doc/extensions/proposed/sycl_ext_oneapi_joint_for.asciidoc
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,261 @@ | ||
= sycl_ext_oneapi_joint_for | ||
|
||
:source-highlighter: coderay | ||
:coderay-linenums-mode: table | ||
|
||
// This section needs to be after the document title. | ||
:doctype: book | ||
:toc2: | ||
:toc: left | ||
:encoding: utf-8 | ||
:lang: en | ||
:dpcpp: pass:[DPC++] | ||
|
||
// Set the default source code type in this document to C++, | ||
// for syntax highlighting purposes. This is needed because | ||
// docbook uses c++ and html5 uses cpp. | ||
:language: {basebackend@docbook:c++:cpp} | ||
|
||
|
||
== Notice | ||
|
||
[%hardbreaks] | ||
Copyright (C) 2024 Intel Corporation. All rights reserved. | ||
|
||
Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks | ||
of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by | ||
permission by Khronos. | ||
|
||
|
||
== Contact | ||
|
||
To report problems with this extension, please open a new issue at: | ||
|
||
https://github.com/intel/llvm/issues | ||
|
||
|
||
== Dependencies | ||
|
||
This extension is written against the SYCL 2020 revision 9 specification. All | ||
references below to the "core SYCL specification" or to section numbers in the | ||
SYCL specification refer to that revision. | ||
|
||
|
||
== Status | ||
|
||
This is a proposed extension specification, intended to gather community | ||
feedback. Interfaces defined in this specification may not be implemented yet | ||
or may be in a preliminary state. The specification itself may also change in | ||
incompatible ways before it is finalized. *Shipping software products should | ||
not rely on APIs defined in this specification.* | ||
|
||
|
||
== Overview | ||
|
||
A very common idiom in SPMD programming sees developers distribute a range over | ||
the members of a group, by writing a loop that looks like this: | ||
|
||
[source,c++] | ||
---- | ||
for (int i = it.get_local_id(0); i < N; i += it.get_local_range(0)) { ... } | ||
---- | ||
|
||
The combination of setting the initial value based on the work-item index and | ||
setting the stride based on the number of work-items in the group ensures that | ||
any data accesses within the loop are appropriately interleaved to maximize | ||
performance. | ||
|
||
Although common, this idiom has several drawbacks: it is difficult to teach to | ||
developers unfamiliar with SPMD programming; it is easy to use incorrectly | ||
(e.g., by forgetting to set a different initial value on each work-item); and | ||
it is difficult to use generically (i.e., for different groups of work-items). | ||
|
||
This extension proposes to encapsulate this pattern in group algorithms, to | ||
simplify generic usage across different group types. | ||
|
||
|
||
== Specification | ||
|
||
=== Feature test macro | ||
|
||
This extension provides a feature-test macro as described in the core SYCL | ||
specification. An implementation supporting this extension must predefine the | ||
macro `SYCL_EXT_ONEAPI_JOINT_FOR` to one of the values defined in the table | ||
below. Applications can test for the existence of this macro to determine if | ||
the implementation supports this feature, or applications can test the macro's | ||
value to determine which of the extension's features the implementation | ||
supports. | ||
|
||
|
||
[%header,cols="1,5"] | ||
|=== | ||
|Value | ||
|Description | ||
|
||
|1 | ||
|The APIs of this experimental extension are not versioned, so the | ||
feature-test macro always has this value. | ||
|=== | ||
|
||
|
||
=== Group algorithms | ||
|
||
[source,c++] | ||
---- | ||
namespace sycl::ext::oneapi::experimental { | ||
|
||
template <typename Group, typename InputIterator, typename Function> | ||
void joint_for_each(Group g, InputIterator first, InputIterator last, Function f); | ||
|
||
} // namespace sycl::ext::oneapi::experimental | ||
---- | ||
|
||
_Constraints_: Available only if `sycl::is_group_v<std::decay_t<Group>>` is | ||
true and `InputIterator` is a random access iterator. | ||
|
||
_Preconditions_: `first` and `last` must be the same for all work-items in | ||
group `g`, and `f` must be an immutable callable with the same type and state | ||
for all work-items in group `g`. | ||
|
||
_Effects_: Blocks until all work-items in group `g` have | ||
reached this synchronization point, then applies `f` to the result of | ||
dereferencing every iterator in the range [`first`, `last`). | ||
The range is distributed across the work-items in group `g`. | ||
|
||
_Synchronization_: The call to this function in each work-item | ||
happens before the algorithm begins execution. | ||
The completion of the algorithm happens before any work-item | ||
blocking on the same synchronization point is unblocked. | ||
|
||
_Remarks_: If `f` returns a result, the result is ignored. | ||
gmlueck marked this conversation as resolved.
Show resolved
Hide resolved
|
||
If the range [`first`, `last`) is not evenly divisible by the number of | ||
work-items in group `g`, the number of times that `f` is invoked will differ | ||
across work-items. | ||
|
||
[source,c++] | ||
---- | ||
namespace sycl::ext::oneapi::experimental { | ||
|
||
template <typename Group, typename Integer, typename Function> | ||
void joint_for(Group g, Integer first, Integer last, Function f); | ||
|
||
} // namespace sycl::ext::oneapi::experimental | ||
---- | ||
|
||
_Constraints_: Available only if `sycl::is_group_v<std::decay_t<Group>>` is | ||
true and `std::is_integral_v<Integer>` is true. | ||
|
||
_Preconditions_: `first` and `last` must be the same for all work-items in | ||
group `g`, and `f` must be an immutable callable with the same type and state | ||
for all work-items in group `g`. | ||
|
||
_Effects_: Blocks until all work-items in group `g` have | ||
reached this synchronization point, then applies `f` to every | ||
integer in the range [`first`, `last`). | ||
The range is distributed across the work-items in group `g`. | ||
|
||
_Synchronization_: The call to this function in each work-item | ||
happens before the algorithm begins execution. | ||
The completion of the algorithm happens before any work-item | ||
blocking on the same synchronization point is unblocked. | ||
|
||
_Remarks_: If `f` returns a result, the result is ignored. | ||
If the range [`first`, `last`) is not evenly divisible by the number of | ||
work-items in group `g`, the number of times that `f` is invoked will differ | ||
across work-items. | ||
|
||
|
||
=== Usage example | ||
|
||
Both `joint_for_each` and `joint_for` can be used to distribute a loop over the | ||
work-items within a group. | ||
|
||
`joint_for_each` can be used to iterate over (and potentially modify) data when | ||
the range can be described using iterators: | ||
|
||
[source,c++] | ||
---- | ||
using syclex = sycl::ext::oneapi::experimental; | ||
|
||
void scale(sycl::queue q, float* data, int N, float factor) { | ||
q.parallel_for(sycl::nd_range<1>(G, L), [=](sycl::nd_item<1> it) { | ||
|
||
// identify the data this group will operate on | ||
auto g = it.get_group(); | ||
int per_group = N / g.get_group_linear_range(); | ||
float* offset = data + per_group * g.get_group_linear_id(); | ||
|
||
// distribute the data over work-items in the group | ||
syclex::joint_for_each(g, offset, offset + per_group, [=](float& x) { | ||
x *= factor; | ||
}); | ||
|
||
}); | ||
} | ||
---- | ||
|
||
`joint_for` can be used when iterators are not available, and is a shortcut for | ||
migrating integer-based `for` loops: | ||
|
||
[source,c++] | ||
---- | ||
using syclex = sycl::ext::oneapi::experimental; | ||
|
||
void scale(sycl::queue q, float* data, int N, float factor) { | ||
q.parallel_for(sycl::nd_range<1>(G, L), [=](sycl::nd_item<1> it) { | ||
|
||
// identify the data this group will operate on | ||
auto g = it.get_group(); | ||
int per_group = N / g.get_group_linear_range(); | ||
float* offset = data + per_group * g.get_group_linear_id(); | ||
|
||
// distribute the data over work-items in the group | ||
syclex::joint_for(g, 0, per_group, [=](int i) { | ||
offset[i] *= factor; | ||
}); | ||
|
||
}); | ||
} | ||
---- | ||
|
||
|
||
== Implementation notes | ||
|
||
This non-normative section provides information about one possible | ||
implementation of this extension. It is not part of the specification of the | ||
extension's API. | ||
|
||
A simple sketch of a possible implementation of this extension, which does not | ||
include the complexities of SFINAE and robust error checking, is given below: | ||
|
||
[source,c++] | ||
---- | ||
template <typename Group, typename InputIterator, typename Function> | ||
void joint_for_each(Group g, InputIterator first, InputIterator last, Function f) { | ||
sycl::group_barrier(g); | ||
typename std::iterator_traits<InputIterator>::difference_type offset = g.get_local_linear_id(); | ||
typename std::iterator_traits<InputIterator>::difference_type stride = g:get_local_linear_range(); | ||
for (InputIterator p = first + offset; p < last; p += stride) { | ||
f(*p); | ||
} | ||
sycl::group_barrier(g); | ||
return f; | ||
} | ||
|
||
template <typename Group, typename Integer, typename Function> | ||
std::enable_if_t<std::is_integral_v<Integer>> | ||
joint_for(Group g, Integer first, Integer last, Function f) { | ||
sycl::group_barrier(g); | ||
Integer offset = g.get_local_linear_id(); | ||
Integer stride = g.get_local_linear_range(); | ||
for (Integer p = first + offset; p < last; p += stride) { | ||
f(p); | ||
} | ||
sycl::group_barrier(g); | ||
return f; | ||
} | ||
---- | ||
|
||
== Issues | ||
|
||
None. |
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Why that and not something like
?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The rest of the functions in the group algorithms library are based on
std::
algorithms, (e.g.,std::reduce
becomessycl::joint_reduce
,std::exclusive_scan
becomessycl::joint_exclusive_scan
) and so following that precedent here seems the most natural:std::for_each
becomessycl::joint_for_each
.