From 93e784e0b60d854194e104c140513493836e235e Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 4 Oct 2024 10:20:22 +0100 Subject: [PATCH 01/15] Add sycl_khr_group_interface extension This extension introduces an alternative interface for groups of work-items, offering several improvements over the SYCL 2020 interface: - Shorter names for member functions, dropping the get_ prefix. - Cleaner separation between properties of a group (e.g., a group id) and properties of the calling work-item (e.g., its id within a group). - Clearer distinction between "group" concept and "work_group" class. - New work_item class to represent a single work-item within a specific parent group. This class also satisfies the group concept, modeling a group containing a single work-item. --- adoc/extensions/index.adoc | 2 + adoc/extensions/sycl_khr_group_interface.adoc | 443 ++++++++++++++++++ 2 files changed, 445 insertions(+) create mode 100644 adoc/extensions/sycl_khr_group_interface.adoc diff --git a/adoc/extensions/index.adoc b/adoc/extensions/index.adoc index 07062df6..7172f9af 100644 --- a/adoc/extensions/index.adoc +++ b/adoc/extensions/index.adoc @@ -11,3 +11,5 @@ specification, but their design is subject to change. // leveloffset=2 allows extensions to be written as standalone documents // include::sycl_khr_extension_name.adoc[leveloffset=2] + +include::sycl_khr_group_interface.adoc[leveloffset=2] diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc new file mode 100644 index 00000000..98c5872d --- /dev/null +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -0,0 +1,443 @@ +[[sec:khr-group-interface]] += SYCL_KHR_GROUP_INTERFACE + +This extension provides an alternative interface for groups of work-items +(including work-groups, sub-groups, and individual work-items) that is simpler +and less verbose than the interface provided by `sycl::group` and +`sycl::sub_group` in SYCL 2020. + +[[sec:khr-group-interface-dependencies]] +== Dependencies + +This extension has no dependencies on other extensions. + +Some features of this extension are only available when an implementation +provides ``. + +[[sec:khr-group-interface-feature-test]] +== Feature test macro + +An implementation supporting this extension must predefine the macro +[code]#SYCL_KHR_GROUP_INTERFACE# to one of the values defined in the table +below. + +[%header,cols="1,5"] +|=== +|Value +|Description + +|1 +|Initial version of this extension. +|=== + +[[sec:khr-group-interface-common]] +== Common group interface + +The `khr::work_group`, `khr::sub_group` and `khr::work_item` objects defined by +this extension all implement a common set of operations. + +[source,role=synopsis] +---- + +namespace sycl { +namespace khr { + +class GroupT { + + public: + using id_type = /* ... */; + using linear_id_type = /* ... */; + using range_type = /* ... */; +#if defined(__cpp_lib_mdspan) + using extents_type = /* ... */; + using index_type = typename extents_type::index_type; // exposition only + using rank_type = typename extents_type::rank_type; // exposition only +#endif + using size_type = /* ... */; + static constexpr int dimensions = /* ... */; + static constexpr memory_scope fence_scope = /* ... */; + + id_type id() const noexcept; + linear_id_type linear_id() const noexcept; + + range_type range() const noexcept; + +#if defined (__cpp_lib_mdspan) + constexpr extents_type extents() const noexcept; + constexpr index_type extent(rank_type r) const noexcept; +#endif + constexpr size_type size() const noexcept; + +}; + +template +work_item get_item(Group g) noexcept; + +template +bool leader_of(Group g) noexcept; + +} // namespace khr +} // namespace sycl +---- + +[[sec:khr-group-interface-common-member-funcs]] +=== Member functions + +.[apidef]#GroupT::id# +[source,role=synopsis,id=api:khr-group-interface-common-group-id] +---- +id_type id() const noexcept +---- + +_Returns_: The index of this group within the index space returned by +[code]#range()#. + +''' + +.[apidef]#GroupT::linear_id# +[source,role=synopsis,id=api:khr-group-interface-common-group-linear-id] +---- +linear_id_type linear_id() const noexcept +---- + +_Returns_: The linearized index (see <>) of this +group within the index space returned by [code]#range()#. + +''' + +.[apidef]#GroupT::range# +[source,role=synopsis,id=api:khr-group-interface-common-group-range] +---- +range_type range() const noexcept +---- + +_Returns_: An index space representing the collection of groups that includes +this group, and which defines the range of valid [code]#id# values for this +group. + +''' + +.[apidef]#GroupT::extents# +[source,role=synopsis,id=api:khr-group-interface-common-group-extents] +---- +constexpr extents_type extents() const noexcept +---- + +_Returns_: The number of work-items in each dimension of the group. + +''' + +.[apidef]#GroupT::extent# +[source,role=synopsis,id=api:khr-group-interface-common-group-extent] +---- +constexpr index_type extent(rank_type r) const noexcept +---- + +_Preconditions_: [code]#r < dimensions# is [code]#true#. + +_Returns_: The number of work-items in the specified dimension of the group. + +''' + +.[apidef]#GroupT::size# +[source,role=synopsis,id=api:common-group-size] +---- +constexpr size_type size() const noexcept +---- + +_Returns_: The total number of work-items in the group, equal to the product of +the number of work-items in each dimension of the group. + +''' + +[[sec:khr-group-interface-common-non-member-funcs]] +==== Non-member functions + +.[apidef]#get_item# +[source,role=synopsis,id=api:common-group-get-item] +---- +template +work_item get_item(Group g) noexcept +---- + +_Returns_: A [code]#work_item# representing the calling work-item within group +[code]#g#. + +''' + +.[apidef]#leader_of# +[source,role=synopsis,id=api:common-group-leader_of] +---- +template +bool leader_of(Group g) noexcept +---- + +_Returns_: [code]#true# if the calling work-item is the leader of group +[code]#g#, and [code]#false# otherwise. + +_Remarks_: [code]#leader_of# returns [code]#true# for only one work-item in a +group. +The leader of the group is determined during construction of the group, and is +invariant for the lifetime of the group. +The leader of the group is guaranteed to be the work-item with index 0 within +the group. + +[[sec:khr-group-interface-work_group]] +== [code]#work_group# class + +The [code]#work_group# class template encapsulates all functionality required to +represent a particular <> within a kernel. +It is not user-constructible. + +The SYCL [code]#work_group# class template provides common by-value semantics +(see <>) and the common group interface (see +<>). + +[source,role=synopsis] +---- +namespace sycl { +namespace khr { + +template +class work_group { + + public: + using id_type = id; + using linear_id_type = size_t; + using range_type = range; +#if defined(__cpp_lib_mdspan) + using extents_type = std::dextents; +#endif + using size_type = size_t; + static constexpr int dimensions = Dimensions; + static constexpr memory_scope fence_scope = memory_scope::work_group; + + work_group(group g) noexcept; + + operator group() const noexcept; + + /* -- common by-value interface members -- */ + + /* -- common group interface members -- */ + +}; + +} // namespace khr +} // namespace sycl +---- + +.[apidef]#work_group constructor# +[source,role=synopsis,id=api:khr-group-interface-work-group-constructor] +---- +work_group(group g) noexcept +---- + +_Effects_: Constructs a [code]#work_group# representing the same collection of +work-items as [code]#g#. + +''' + +.[apidef]#work_group conversion operator# +[source,role=synopsis,id=api:khr-group-interface-work-group-conversion-operator] +---- +operator group() const noexcept; +---- + +_Returns_: A [code]#group# representing the same collection of work-items as +this [code]#work_group#. + +''' + +[[sec:khr-group-interface-sub_group]] +== [code]#sub_group# class + +The [code]#sub_group# class template encapsulates all functionality required to +represent a particular <> within a kernel. +It is not user-constructible. + +The SYCL [code]#sub_group# class template provides common by-value semantics +(see <>) and the common group interface (see +<>). + +[source,role=synopsis] +---- +namespace sycl { +namespace khr { + +class sub_group { + + public: + using id_type = id<1>; + using linear_id_type = uint32_t; + using range_type = range<1>; +#if defined(__cpp_lib_mdspan) + using extents_type = std::dextents; +#endif + using size_type = uint32_t; + static constexpr int dimensions = 1; + static constexpr memory_scope fence_scope = memory_scope::sub_group; + + sub_group(sycl::sub_group sg) noexcept; + + operator sycl::sub_group() const noexcept; + + constexpr size_type max_size() const noexcept; + + /* -- common by-value interface members -- */ + + /* -- common group interface members -- */ + +}; + +} // namespace khr +} // namespace sycl +---- + +.[apidef]#sub_group constructor# +[source,role=synopsis,id=api:khr-group-interface-sub-group-constructor] +---- +sub_group(sycl::sub_group sg) noexcept +---- + +_Effects_: Constructs a [code]#sub_group# representing the same collection of +work-items as [code]#sg#. + +''' + +.[apidef]#sub_group conversion operator# +[source,role=synopsis,id=api:khr-group-interface-sub-group-conversion-operator] +---- +operator sycl::sub_group() const noexcept; +---- + +_Returns_: A [code]#sycl::sub_group# representing the same collection of +work-items as this [code]#sub_group#. + +''' + +.[apidef]#max_size# +[source,role=synopsis,id=api:khr-group-interface-sub-group-max-size] +---- +constexpr size_type max_size() const noexcept; +---- + +_Returns_: The maximum number of work-items permitted in any <> for +the executing kernel. + +{note}There is no guarantee that any sub-group within the work-group contains +the maximum number of work-items.{endnote} + +_Remarks_: The value returned by this function must reflect the value passed to +the [code]#reqd_sub_group_size# attribute, if present. +If no such attribute is present, the value returned is determined by the +<>. + +''' + +[[sec:khr-group-interface-work_item]] +== [code]#work_item# class + +The [code]#work_item# class template encapsulates all functionality required to +represent a single <> within a kernel. +It is not user-constructible. + +The SYCL [code]#work_item# class template provides common by-value semantics +(see <>) and the common group interface (see +<>). + +[source,role=synopsis] +---- +namespace sycl { +namespace khr { + +template +class work_item { + + public: + using id_type = typename ParentGroup::id_type; + using linear_id_type = typename ParentGroup::linear_id_type; + using range_type = typename ParentGroup::range_type; +#if defined(__cpp_lib_mdspan) + using extents_type = std::extents; +#endif + using size_type = typename ParentGroup::size_type; + static constexpr int dimensions = ParentGroup::dimensions; + static constexpr memory_scope fence_scope = memory_scope::work_item; + + /* -- common by-value interface members -- */ + + /* -- common group interface members -- */ + +}; + +} // namespace khr +} // namespace sycl +---- + +[[sec:khr-group-interface-example]] +== Example + +The example below demonstrates the usage of this extension. + +[source,,linenums] +---- +#include +#include +#include +#include +using namespace sycl; // (optional) avoids need for "sycl::" before SYCL name + +constexpr size_t N = 1024; +constexpr size_t M = 256; + +int main() { + + queue q; + + int* in = malloc_shared(N * M, q); + int* out = malloc_shared(N, q); + + std::iota(in, in + N * M, 0); + std::fill(out, out + N, 0); + + q.parallel_for(nd_range<1>{64, 32}, [=](nd_item<1> ndit) { + + // opt into the new group interface + khr::work_group<1> g = ndit.get_group(); + khr::work_item<1> it = get_item(g); + + // distribute N loop over work-groups + for (size_t i = g.linear_id(); i < N; i += g.range().size()) { + + // distribute M loop over work-items in the work-group + int sum = 0; + for (size_t j = it.linear_id(); j < M; j += it.range().size()) { + sum += in[i * M + j]; + } + + // accumulate partial results and write out + sum = sycl::reduce_over_group((sycl::group<1>) g, sum, sycl::plus<>()); + if (khr::leader_of(g)) { + out[i] = sum; + } + + } + + }).wait(); + + std::cout << std::endl << "Result:" << std::endl; + for (size_t i = 0; i < N; i++) { + int sum = 0; + for (size_t j = 0; j < M; j++) { + sum += in[i * M + j]; + } + if (sum != out[i]) { + std::cout << "Wrong value " << out[i] << " on element " << i << std::endl; + exit(-1); + } + } + + std::cout << "Good computation!" << std::endl; + return 0; +} +---- From eb009d06807c4bb2c58ce8fd3051f779894e87d7 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 4 Oct 2024 13:27:07 +0100 Subject: [PATCH 02/15] Fix typo in khr_group_interface example khr::work_item is not templated on dimensions. --- adoc/extensions/sycl_khr_group_interface.adoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 98c5872d..c07298a3 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -404,7 +404,7 @@ int main() { // opt into the new group interface khr::work_group<1> g = ndit.get_group(); - khr::work_item<1> it = get_item(g); + khr::work_item it = get_item(g); // distribute N loop over work-groups for (size_t i = g.linear_id(); i < N; i += g.range().size()) { From 04bc66d513052c4fa8993593041e0f67cae548af Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 4 Oct 2024 16:47:54 +0100 Subject: [PATCH 03/15] Fix khr::work_item extents_type The synopsis previously assumed that the extents would be one dimensional, but it must match the ParentGroup. There is no suitable shorthand alias or exposition-only description for this case defined by mdspan, so replace the type with a comment. --- adoc/extensions/sycl_khr_group_interface.adoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index c07298a3..57c30390 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -358,7 +358,7 @@ class work_item { using linear_id_type = typename ParentGroup::linear_id_type; using range_type = typename ParentGroup::range_type; #if defined(__cpp_lib_mdspan) - using extents_type = std::extents; + using extents_type = /* extents of all 1s with ParentGroup's index type */ #endif using size_type = typename ParentGroup::size_type; static constexpr int dimensions = ParentGroup::dimensions; From d6e0d11d9456985bf96beb75ccd60e0a821762c6 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Mon, 7 Oct 2024 08:49:00 +0100 Subject: [PATCH 04/15] Add constraints to get_item --- adoc/extensions/sycl_khr_group_interface.adoc | 2 ++ 1 file changed, 2 insertions(+) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 57c30390..27c1561b 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -160,6 +160,8 @@ template work_item get_item(Group g) noexcept ---- +_Constraints_: Available only if `Group` is `work_group` or `sub_group`. + _Returns_: A [code]#work_item# representing the calling work-item within group [code]#g#. From 008ac23565df212244a692bd01711766f8433000 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 5 Dec 2024 10:47:21 +0000 Subject: [PATCH 05/15] Add extent observers to common group interface rank(), rank_dynamic() and static_extent() are all exposed as static member functions in mdspan, so we follow that precedent here. --- adoc/extensions/sycl_khr_group_interface.adoc | 34 +++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 27c1561b..942ef514 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -65,6 +65,10 @@ class GroupT { #if defined (__cpp_lib_mdspan) constexpr extents_type extents() const noexcept; constexpr index_type extent(rank_type r) const noexcept; + + static constexpr rank_type rank() noexcept; + static constexpr rank_type rank_dynamic() noexcept; + static constexpr size_t static_extent(rank_type r) noexcept; #endif constexpr size_type size() const noexcept; @@ -139,6 +143,36 @@ _Returns_: The number of work-items in the specified dimension of the group. ''' +.[apidef]#GroupT::rank# +[source,role=synopsis,id=api:khr-group-interface-common-group-rank] +---- +static constexpr rank_type rank() noexcept; +---- + +_Effects_: Equivalent to [code]#return extents_type::rank();#. + +''' + +.[apidef]#GroupT::rank_dynamic# +[source,role=synopsis,id=api:khr-group-interface-common-group-rank_dynamic] +---- +static constexpr rank_type rank_dynamic() noexcept; +---- + +_Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. + +''' + +.[apidef]#GroupT::static_extent# +[source,role=synopsis,id=api:khr-group-interface-common-group-static_extent] +---- +static constexpr size_t static_extent(rank_type r) noexcept; +---- + +_Effects_: Equivalent to [code]#return extents_type::static_extent(r);#. + +''' + .[apidef]#GroupT::size# [source,role=synopsis,id=api:common-group-size] ---- From d59ccb99b49f80a5e7084daf719cb8713877d76c Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 5 Dec 2024 11:00:32 +0000 Subject: [PATCH 06/15] Add semicolons to khr_group_interface synopses --- adoc/extensions/sycl_khr_group_interface.adoc | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 942ef514..8c96397b 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -90,7 +90,7 @@ bool leader_of(Group g) noexcept; .[apidef]#GroupT::id# [source,role=synopsis,id=api:khr-group-interface-common-group-id] ---- -id_type id() const noexcept +id_type id() const noexcept; ---- _Returns_: The index of this group within the index space returned by @@ -101,7 +101,7 @@ _Returns_: The index of this group within the index space returned by .[apidef]#GroupT::linear_id# [source,role=synopsis,id=api:khr-group-interface-common-group-linear-id] ---- -linear_id_type linear_id() const noexcept +linear_id_type linear_id() const noexcept; ---- _Returns_: The linearized index (see <>) of this @@ -112,7 +112,7 @@ group within the index space returned by [code]#range()#. .[apidef]#GroupT::range# [source,role=synopsis,id=api:khr-group-interface-common-group-range] ---- -range_type range() const noexcept +range_type range() const noexcept; ---- _Returns_: An index space representing the collection of groups that includes @@ -124,7 +124,7 @@ group. .[apidef]#GroupT::extents# [source,role=synopsis,id=api:khr-group-interface-common-group-extents] ---- -constexpr extents_type extents() const noexcept +constexpr extents_type extents() const noexcept; ---- _Returns_: The number of work-items in each dimension of the group. @@ -134,7 +134,7 @@ _Returns_: The number of work-items in each dimension of the group. .[apidef]#GroupT::extent# [source,role=synopsis,id=api:khr-group-interface-common-group-extent] ---- -constexpr index_type extent(rank_type r) const noexcept +constexpr index_type extent(rank_type r) const noexcept; ---- _Preconditions_: [code]#r < dimensions# is [code]#true#. @@ -176,7 +176,7 @@ _Effects_: Equivalent to [code]#return extents_type::static_extent(r);#. .[apidef]#GroupT::size# [source,role=synopsis,id=api:common-group-size] ---- -constexpr size_type size() const noexcept +constexpr size_type size() const noexcept; ---- _Returns_: The total number of work-items in the group, equal to the product of @@ -191,7 +191,7 @@ the number of work-items in each dimension of the group. [source,role=synopsis,id=api:common-group-get-item] ---- template -work_item get_item(Group g) noexcept +work_item get_item(Group g) noexcept; ---- _Constraints_: Available only if `Group` is `work_group` or `sub_group`. @@ -205,7 +205,7 @@ _Returns_: A [code]#work_item# representing the calling work-item within group [source,role=synopsis,id=api:common-group-leader_of] ---- template -bool leader_of(Group g) noexcept +bool leader_of(Group g) noexcept; ---- _Returns_: [code]#true# if the calling work-item is the leader of group @@ -265,7 +265,7 @@ class work_group { .[apidef]#work_group constructor# [source,role=synopsis,id=api:khr-group-interface-work-group-constructor] ---- -work_group(group g) noexcept +work_group(group g) noexcept; ---- _Effects_: Constructs a [code]#work_group# representing the same collection of @@ -332,7 +332,7 @@ class sub_group { .[apidef]#sub_group constructor# [source,role=synopsis,id=api:khr-group-interface-sub-group-constructor] ---- -sub_group(sycl::sub_group sg) noexcept +sub_group(sycl::sub_group sg) noexcept; ---- _Effects_: Constructs a [code]#sub_group# representing the same collection of From a1a7e8c675cc303be7d565a1a06d15e89f0685c7 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 10 Dec 2024 09:28:18 +0000 Subject: [PATCH 07/15] Rename GroupT to __group__ The name __group__ is consistent with the names of the swizzle classes. --- adoc/extensions/sycl_khr_group_interface.adoc | 27 +++++++++++-------- 1 file changed, 16 insertions(+), 11 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 8c96397b..2645d83f 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -34,7 +34,12 @@ below. == Common group interface The `khr::work_group`, `khr::sub_group` and `khr::work_item` objects defined by -this extension all implement a common set of operations. +this extension all implement a common set of operations, which are shown in the +synopsis below. +The name [code]#__group__# in that synopsis is a placeholder for each of these +three types. +When the synopsis shows an ellipsis ([code]#+/*...*/+#), the subsequent sections +clarify the definition for each type. [source,role=synopsis] ---- @@ -42,7 +47,7 @@ this extension all implement a common set of operations. namespace sycl { namespace khr { -class GroupT { +class __group__ { public: using id_type = /* ... */; @@ -87,7 +92,7 @@ bool leader_of(Group g) noexcept; [[sec:khr-group-interface-common-member-funcs]] === Member functions -.[apidef]#GroupT::id# +.[apidef]#__group__::id# [source,role=synopsis,id=api:khr-group-interface-common-group-id] ---- id_type id() const noexcept; @@ -98,7 +103,7 @@ _Returns_: The index of this group within the index space returned by ''' -.[apidef]#GroupT::linear_id# +.[apidef]#__group__::linear_id# [source,role=synopsis,id=api:khr-group-interface-common-group-linear-id] ---- linear_id_type linear_id() const noexcept; @@ -109,7 +114,7 @@ group within the index space returned by [code]#range()#. ''' -.[apidef]#GroupT::range# +.[apidef]#__group__::range# [source,role=synopsis,id=api:khr-group-interface-common-group-range] ---- range_type range() const noexcept; @@ -121,7 +126,7 @@ group. ''' -.[apidef]#GroupT::extents# +.[apidef]#__group__::extents# [source,role=synopsis,id=api:khr-group-interface-common-group-extents] ---- constexpr extents_type extents() const noexcept; @@ -131,7 +136,7 @@ _Returns_: The number of work-items in each dimension of the group. ''' -.[apidef]#GroupT::extent# +.[apidef]#__group__::extent# [source,role=synopsis,id=api:khr-group-interface-common-group-extent] ---- constexpr index_type extent(rank_type r) const noexcept; @@ -143,7 +148,7 @@ _Returns_: The number of work-items in the specified dimension of the group. ''' -.[apidef]#GroupT::rank# +.[apidef]#__group__::rank# [source,role=synopsis,id=api:khr-group-interface-common-group-rank] ---- static constexpr rank_type rank() noexcept; @@ -153,7 +158,7 @@ _Effects_: Equivalent to [code]#return extents_type::rank();#. ''' -.[apidef]#GroupT::rank_dynamic# +.[apidef]#__group__::rank_dynamic# [source,role=synopsis,id=api:khr-group-interface-common-group-rank_dynamic] ---- static constexpr rank_type rank_dynamic() noexcept; @@ -163,7 +168,7 @@ _Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. ''' -.[apidef]#GroupT::static_extent# +.[apidef]#__group__::static_extent# [source,role=synopsis,id=api:khr-group-interface-common-group-static_extent] ---- static constexpr size_t static_extent(rank_type r) noexcept; @@ -173,7 +178,7 @@ _Effects_: Equivalent to [code]#return extents_type::static_extent(r);#. ''' -.[apidef]#GroupT::size# +.[apidef]#__group__::size# [source,role=synopsis,id=api:common-group-size] ---- constexpr size_type size() const noexcept; From dd2110a6f71f9ca00d688296dc6f7f3dcaf0bfb8 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 10 Dec 2024 09:36:34 +0000 Subject: [PATCH 08/15] Remove exposition-only types from group interface --- adoc/extensions/sycl_khr_group_interface.adoc | 16 +++++++--------- 1 file changed, 7 insertions(+), 9 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 2645d83f..01a43570 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -55,8 +55,6 @@ class __group__ { using range_type = /* ... */; #if defined(__cpp_lib_mdspan) using extents_type = /* ... */; - using index_type = typename extents_type::index_type; // exposition only - using rank_type = typename extents_type::rank_type; // exposition only #endif using size_type = /* ... */; static constexpr int dimensions = /* ... */; @@ -69,10 +67,10 @@ class __group__ { #if defined (__cpp_lib_mdspan) constexpr extents_type extents() const noexcept; - constexpr index_type extent(rank_type r) const noexcept; + constexpr extents_type::index_type extent(extents_type::rank_type r) const noexcept; - static constexpr rank_type rank() noexcept; - static constexpr rank_type rank_dynamic() noexcept; + static constexpr extents_type::rank_type rank() noexcept; + static constexpr extents_type::rank_type rank_dynamic() noexcept; static constexpr size_t static_extent(rank_type r) noexcept; #endif constexpr size_type size() const noexcept; @@ -139,7 +137,7 @@ _Returns_: The number of work-items in each dimension of the group. .[apidef]#__group__::extent# [source,role=synopsis,id=api:khr-group-interface-common-group-extent] ---- -constexpr index_type extent(rank_type r) const noexcept; +constexpr extents_type::index_type extent(extents_type::rank_type r) const noexcept; ---- _Preconditions_: [code]#r < dimensions# is [code]#true#. @@ -151,7 +149,7 @@ _Returns_: The number of work-items in the specified dimension of the group. .[apidef]#__group__::rank# [source,role=synopsis,id=api:khr-group-interface-common-group-rank] ---- -static constexpr rank_type rank() noexcept; +static constexpr extents_type::rank_type rank() noexcept; ---- _Effects_: Equivalent to [code]#return extents_type::rank();#. @@ -161,7 +159,7 @@ _Effects_: Equivalent to [code]#return extents_type::rank();#. .[apidef]#__group__::rank_dynamic# [source,role=synopsis,id=api:khr-group-interface-common-group-rank_dynamic] ---- -static constexpr rank_type rank_dynamic() noexcept; +static constexpr extents_type::rank_type rank_dynamic() noexcept; ---- _Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. @@ -171,7 +169,7 @@ _Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. .[apidef]#__group__::static_extent# [source,role=synopsis,id=api:khr-group-interface-common-group-static_extent] ---- -static constexpr size_t static_extent(rank_type r) noexcept; +static constexpr size_t static_extent(extents_type::rank_type r) noexcept; ---- _Effects_: Equivalent to [code]#return extents_type::static_extent(r);#. From da495737245c552706be1450c5c127bb645ae81d Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 10 Dec 2024 09:39:06 +0000 Subject: [PATCH 09/15] Collapse sycl::khr namespace in group interface --- adoc/extensions/sycl_khr_group_interface.adoc | 24 +++++++------------ 1 file changed, 8 insertions(+), 16 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 01a43570..b98a1668 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -44,8 +44,7 @@ clarify the definition for each type. [source,role=synopsis] ---- -namespace sycl { -namespace khr { +namespace sycl::khr { class __group__ { @@ -83,8 +82,7 @@ work_item get_item(Group g) noexcept; template bool leader_of(Group g) noexcept; -} // namespace khr -} // namespace sycl +} // namespace sycl::khr ---- [[sec:khr-group-interface-common-member-funcs]] @@ -234,8 +232,7 @@ The SYCL [code]#work_group# class template provides common by-value semantics [source,role=synopsis] ---- -namespace sycl { -namespace khr { +namespace sycl::khr { template class work_group { @@ -261,8 +258,7 @@ class work_group { }; -} // namespace khr -} // namespace sycl +} // namespace sycl::khr ---- .[apidef]#work_group constructor# @@ -300,8 +296,7 @@ The SYCL [code]#sub_group# class template provides common by-value semantics [source,role=synopsis] ---- -namespace sycl { -namespace khr { +namespace sycl::khr { class sub_group { @@ -328,8 +323,7 @@ class sub_group { }; -} // namespace khr -} // namespace sycl +} // namespace sycl::khr ---- .[apidef]#sub_group constructor# @@ -386,8 +380,7 @@ The SYCL [code]#work_item# class template provides common by-value semantics [source,role=synopsis] ---- -namespace sycl { -namespace khr { +namespace sycl::khr { template class work_item { @@ -409,8 +402,7 @@ class work_item { }; -} // namespace khr -} // namespace sycl +} // namespace sycl::khr ---- [[sec:khr-group-interface-example]] From 64e45c5c5947649fa66711bf1c3a9c30e6bd0895 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 10 Dec 2024 09:44:08 +0000 Subject: [PATCH 10/15] Update group interface constraints for consistency --- adoc/extensions/sycl_khr_group_interface.adoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index b98a1668..b74e9175 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -195,7 +195,7 @@ template work_item get_item(Group g) noexcept; ---- -_Constraints_: Available only if `Group` is `work_group` or `sub_group`. +_Constraints_: `Group` is `work_group` or `sub_group`. _Returns_: A [code]#work_item# representing the calling work-item within group [code]#g#. From 80914ee65f942fc98b599c21d1d9f31fd87f65b6 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 10 Dec 2024 09:49:38 +0000 Subject: [PATCH 11/15] Clarify that is C++23 Co-authored-by: Greg Lueck --- adoc/extensions/sycl_khr_group_interface.adoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index b74e9175..f26b5e04 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -12,7 +12,7 @@ and less verbose than the interface provided by `sycl::group` and This extension has no dependencies on other extensions. Some features of this extension are only available when an implementation -provides ``. +provides ``, which is defined in C++23. [[sec:khr-group-interface-feature-test]] == Feature test macro From de356628056c3d1034fcde2e2be59f9c4ebf3ecc Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Tue, 10 Dec 2024 10:01:56 +0000 Subject: [PATCH 12/15] Fix [api], [apidef], etc in group interface --- adoc/extensions/sycl_khr_group_interface.adoc | 38 +++++++++---------- 1 file changed, 19 insertions(+), 19 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index f26b5e04..ca2044e0 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -36,7 +36,7 @@ below. The `khr::work_group`, `khr::sub_group` and `khr::work_item` objects defined by this extension all implement a common set of operations, which are shown in the synopsis below. -The name [code]#__group__# in that synopsis is a placeholder for each of these +The name [code]#+__group__+# in that synopsis is a placeholder for each of these three types. When the synopsis shows an ellipsis ([code]#+/*...*/+#), the subsequent sections clarify the definition for each type. @@ -88,29 +88,29 @@ bool leader_of(Group g) noexcept; [[sec:khr-group-interface-common-member-funcs]] === Member functions -.[apidef]#__group__::id# +.[apidef]#+__group__::id+# [source,role=synopsis,id=api:khr-group-interface-common-group-id] ---- id_type id() const noexcept; ---- _Returns_: The index of this group within the index space returned by -[code]#range()#. +[api]#+__group__::range+#. ''' -.[apidef]#__group__::linear_id# +.[apidef]#+__group__::linear_id+# [source,role=synopsis,id=api:khr-group-interface-common-group-linear-id] ---- linear_id_type linear_id() const noexcept; ---- _Returns_: The linearized index (see <>) of this -group within the index space returned by [code]#range()#. +group within the index space returned by [api]#+__group__::range+#. ''' -.[apidef]#__group__::range# +.[apidef]#+__group__::range+# [source,role=synopsis,id=api:khr-group-interface-common-group-range] ---- range_type range() const noexcept; @@ -122,7 +122,7 @@ group. ''' -.[apidef]#__group__::extents# +.[apidef]#+__group__::extents+# [source,role=synopsis,id=api:khr-group-interface-common-group-extents] ---- constexpr extents_type extents() const noexcept; @@ -132,7 +132,7 @@ _Returns_: The number of work-items in each dimension of the group. ''' -.[apidef]#__group__::extent# +.[apidef]#+__group__::extent+# [source,role=synopsis,id=api:khr-group-interface-common-group-extent] ---- constexpr extents_type::index_type extent(extents_type::rank_type r) const noexcept; @@ -144,7 +144,7 @@ _Returns_: The number of work-items in the specified dimension of the group. ''' -.[apidef]#__group__::rank# +.[apidef]#+__group__::rank+# [source,role=synopsis,id=api:khr-group-interface-common-group-rank] ---- static constexpr extents_type::rank_type rank() noexcept; @@ -154,7 +154,7 @@ _Effects_: Equivalent to [code]#return extents_type::rank();#. ''' -.[apidef]#__group__::rank_dynamic# +.[apidef]#+__group__::rank_dynamic+# [source,role=synopsis,id=api:khr-group-interface-common-group-rank_dynamic] ---- static constexpr extents_type::rank_type rank_dynamic() noexcept; @@ -164,7 +164,7 @@ _Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. ''' -.[apidef]#__group__::static_extent# +.[apidef]#+__group__::static_extent+# [source,role=synopsis,id=api:khr-group-interface-common-group-static_extent] ---- static constexpr size_t static_extent(extents_type::rank_type r) noexcept; @@ -174,7 +174,7 @@ _Effects_: Equivalent to [code]#return extents_type::static_extent(r);#. ''' -.[apidef]#__group__::size# +.[apidef]#+__group__::size+# [source,role=synopsis,id=api:common-group-size] ---- constexpr size_type size() const noexcept; @@ -188,7 +188,7 @@ the number of work-items in each dimension of the group. [[sec:khr-group-interface-common-non-member-funcs]] ==== Non-member functions -.[apidef]#get_item# +.[apidef]#khr::get_item# [source,role=synopsis,id=api:common-group-get-item] ---- template @@ -202,7 +202,7 @@ _Returns_: A [code]#work_item# representing the calling work-item within group ''' -.[apidef]#leader_of# +.[apidef]#khr::leader_of# [source,role=synopsis,id=api:common-group-leader_of] ---- template @@ -261,7 +261,7 @@ class work_group { } // namespace sycl::khr ---- -.[apidef]#work_group constructor# +.[apititle]#work_group constructor# [source,role=synopsis,id=api:khr-group-interface-work-group-constructor] ---- work_group(group g) noexcept; @@ -272,7 +272,7 @@ work-items as [code]#g#. ''' -.[apidef]#work_group conversion operator# +.[apititle]#work_group conversion operator# [source,role=synopsis,id=api:khr-group-interface-work-group-conversion-operator] ---- operator group() const noexcept; @@ -326,7 +326,7 @@ class sub_group { } // namespace sycl::khr ---- -.[apidef]#sub_group constructor# +.[apititle]#sub_group constructor# [source,role=synopsis,id=api:khr-group-interface-sub-group-constructor] ---- sub_group(sycl::sub_group sg) noexcept; @@ -337,7 +337,7 @@ work-items as [code]#sg#. ''' -.[apidef]#sub_group conversion operator# +.[apititle]#sub_group conversion operator# [source,role=synopsis,id=api:khr-group-interface-sub-group-conversion-operator] ---- operator sycl::sub_group() const noexcept; @@ -348,7 +348,7 @@ work-items as this [code]#sub_group#. ''' -.[apidef]#max_size# +.[apidef]#khr::sub_group::max_size# [source,role=synopsis,id=api:khr-group-interface-sub-group-max-size] ---- constexpr size_type max_size() const noexcept; From 3a6e6e2913b96b8b6674dbabb303e38230b27a69 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 12 Dec 2024 10:03:15 +0000 Subject: [PATCH 13/15] Clarify constructability for khr:: group types The previous "not user constructible" wording was based on the description of existing group types. The new wording explains why khr::group, khr::sub_group, and khr::work_item have limited constructors and also points to functions that can be used to acquire instances of these types. --- adoc/extensions/sycl_khr_group_interface.adoc | 34 +++++++++++++++---- 1 file changed, 27 insertions(+), 7 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index ca2044e0..5e70da7a 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -1,4 +1,4 @@ -[[sec:khr-group-interface]] +[sec:khr-group-interface]] = SYCL_KHR_GROUP_INTERFACE This extension provides an alternative interface for groups of work-items @@ -223,8 +223,15 @@ the group. == [code]#work_group# class The [code]#work_group# class template encapsulates all functionality required to -represent a particular <> within a kernel. -It is not user-constructible. +represent a specific <> within a kernel. + +The set of work-items represented by an instance of the [code]#work_group# class +template is determined by the implementation, and there is subsequently no way +for a user to construct arbitrary instances of the [code]#work_group# class +template. +Instances of the [code]#work_group# class template can only be acquired from a +call to a standard SYCL function, or by converting an instance of the +[code]#sycl::group# class template. The SYCL [code]#work_group# class template provides common by-value semantics (see <>) and the common group interface (see @@ -287,8 +294,15 @@ this [code]#work_group#. == [code]#sub_group# class The [code]#sub_group# class template encapsulates all functionality required to -represent a particular <> within a kernel. -It is not user-constructible. +represent a specific <> within a <>. + +The set of work-items represented by an instance of the [code]#sub_group# class +template is determined by the implementation, and there is subsequently no way +for a user to construct arbitrary instances of the [code]#sub_group# class +template. +Instances of the [code]#sub_group# class template can only be acquired from a +call to a standard SYCL function, or by converting an instance of the +[code]#sycl::sub_group# class template. The SYCL [code]#sub_group# class template provides common by-value semantics (see <>) and the common group interface (see @@ -371,8 +385,14 @@ If no such attribute is present, the value returned is determined by the == [code]#work_item# class The [code]#work_item# class template encapsulates all functionality required to -represent a single <> within a kernel. -It is not user-constructible. +represent a single <> within a specific <> of work-items. + +The mechanism used to determine the calling work-item's position within a given +group of work-items is implementation-defined, and there is subsequently no way +for a user to construct arbitrary instances of the [code]#work_item# class +template. +Instances of the [code]#work_item# class template can only be acquired from a +call to [api]#khr::get_item#. The SYCL [code]#work_item# class template provides common by-value semantics (see <>) and the common group interface (see From 9355578a052cee3d71288f241f160ae1f6b55c6a Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 12 Dec 2024 10:13:38 +0000 Subject: [PATCH 14/15] Adopt new standard wording for C++23 dependency --- adoc/extensions/sycl_khr_group_interface.adoc | 43 ++++++++++--------- adoc/syclbase.adoc | 1 + 2 files changed, 23 insertions(+), 21 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 5e70da7a..8ee8ae72 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -11,8 +11,8 @@ and less verbose than the interface provided by `sycl::group` and This extension has no dependencies on other extensions. -Some features of this extension are only available when an implementation -provides ``, which is defined in C++23. +Some features of this extension are only available when a SYCL implementation +conforms to {cpp23} or later. [[sec:khr-group-interface-feature-test]] == Feature test macro @@ -52,9 +52,7 @@ class __group__ { using id_type = /* ... */; using linear_id_type = /* ... */; using range_type = /* ... */; -#if defined(__cpp_lib_mdspan) - using extents_type = /* ... */; -#endif + using extents_type = /* ... */; // C++23 using size_type = /* ... */; static constexpr int dimensions = /* ... */; static constexpr memory_scope fence_scope = /* ... */; @@ -64,14 +62,13 @@ class __group__ { range_type range() const noexcept; -#if defined (__cpp_lib_mdspan) - constexpr extents_type extents() const noexcept; - constexpr extents_type::index_type extent(extents_type::rank_type r) const noexcept; + constexpr extents_type extents() const noexcept; // C++23 + constexpr extents_type::index_type extent(extents_type::rank_type r) const noexcept; // C++23 + + static constexpr extents_type::rank_type rank() noexcept; // C++23 + static constexpr extents_type::rank_type rank_dynamic() noexcept; // C++23 + static constexpr size_t static_extent(rank_type r) noexcept; // C++23 - static constexpr extents_type::rank_type rank() noexcept; - static constexpr extents_type::rank_type rank_dynamic() noexcept; - static constexpr size_t static_extent(rank_type r) noexcept; -#endif constexpr size_type size() const noexcept; }; @@ -128,6 +125,8 @@ group. constexpr extents_type extents() const noexcept; ---- +_Minimum C++ Version_: {cpp23} + _Returns_: The number of work-items in each dimension of the group. ''' @@ -138,6 +137,8 @@ _Returns_: The number of work-items in each dimension of the group. constexpr extents_type::index_type extent(extents_type::rank_type r) const noexcept; ---- +_Minimum C++ Version_: {cpp23} + _Preconditions_: [code]#r < dimensions# is [code]#true#. _Returns_: The number of work-items in the specified dimension of the group. @@ -150,6 +151,8 @@ _Returns_: The number of work-items in the specified dimension of the group. static constexpr extents_type::rank_type rank() noexcept; ---- +_Minimum C++ Version_: {cpp23} + _Effects_: Equivalent to [code]#return extents_type::rank();#. ''' @@ -160,6 +163,8 @@ _Effects_: Equivalent to [code]#return extents_type::rank();#. static constexpr extents_type::rank_type rank_dynamic() noexcept; ---- +_Minimum C++ Version_: {cpp23} + _Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. ''' @@ -170,6 +175,8 @@ _Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#. static constexpr size_t static_extent(extents_type::rank_type r) noexcept; ---- +_Minimum C++ Version_: {cpp23} + _Effects_: Equivalent to [code]#return extents_type::static_extent(r);#. ''' @@ -248,9 +255,7 @@ class work_group { using id_type = id; using linear_id_type = size_t; using range_type = range; -#if defined(__cpp_lib_mdspan) - using extents_type = std::dextents; -#endif + using extents_type = std::dextents; // C++23 using size_type = size_t; static constexpr int dimensions = Dimensions; static constexpr memory_scope fence_scope = memory_scope::work_group; @@ -318,9 +323,7 @@ class sub_group { using id_type = id<1>; using linear_id_type = uint32_t; using range_type = range<1>; -#if defined(__cpp_lib_mdspan) - using extents_type = std::dextents; -#endif + using extents_type = std::dextents; // C++23 using size_type = uint32_t; static constexpr int dimensions = 1; static constexpr memory_scope fence_scope = memory_scope::sub_group; @@ -409,9 +412,7 @@ class work_item { using id_type = typename ParentGroup::id_type; using linear_id_type = typename ParentGroup::linear_id_type; using range_type = typename ParentGroup::range_type; -#if defined(__cpp_lib_mdspan) - using extents_type = /* extents of all 1s with ParentGroup's index type */ -#endif + using extents_type = /* extents of all 1s with ParentGroup's index type */; // C++23 using size_type = typename ParentGroup::size_type; static constexpr int dimensions = ParentGroup::dimensions; static constexpr memory_scope fence_scope = memory_scope::work_item; diff --git a/adoc/syclbase.adoc b/adoc/syclbase.adoc index f6f03041..9a4832bf 100644 --- a/adoc/syclbase.adoc +++ b/adoc/syclbase.adoc @@ -78,6 +78,7 @@ include::config/api_xrefs.adoc[] // Asciidoctor uses ++ to denote spans, so use these attributes instead :cpp17: pass:[C++17] :cpp20: pass:[C++20] +:cpp23: pass:[C++23] // Fortunately, the cpp macro for C++ is defined by default in AsciiDoctor // Use these to format a non-normative note. The Asciidoc source: From b05db85e72f489ea10d8b87f111624719425614d Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Thu, 12 Dec 2024 10:16:18 +0000 Subject: [PATCH 15/15] Fix asciidoc formatting in group_interface --- adoc/extensions/sycl_khr_group_interface.adoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/adoc/extensions/sycl_khr_group_interface.adoc b/adoc/extensions/sycl_khr_group_interface.adoc index 8ee8ae72..4dc6ac15 100644 --- a/adoc/extensions/sycl_khr_group_interface.adoc +++ b/adoc/extensions/sycl_khr_group_interface.adoc @@ -1,4 +1,4 @@ -[sec:khr-group-interface]] +[[sec:khr-group-interface]] = SYCL_KHR_GROUP_INTERFACE This extension provides an alternative interface for groups of work-items @@ -193,7 +193,7 @@ the number of work-items in each dimension of the group. ''' [[sec:khr-group-interface-common-non-member-funcs]] -==== Non-member functions +=== Non-member functions .[apidef]#khr::get_item# [source,role=synopsis,id=api:common-group-get-item]