[SYCL][Doc] Add draft of sycl_ext_oneapi_joint_for (#14886)

Adds a proposed extension for encapsulating the pattern of distributing
loops across the work-items in a group.

---------

Signed-off-by: John Pennycook <[email protected]>
Co-authored-by: Greg Lueck <[email protected]>

Author: Pennycook

sycl/doc/extensions/proposed/sycl_ext_oneapi_joint_for.asciidoc

@@ -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.
+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.