From 93e784e0b60d854194e104c140513493836e235e Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 4 Oct 2024 10:20:22 +0100 Subject: [PATCH] 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; +} +----