diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc index 0ef2126d86fa1..44eb1b376d113 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc @@ -182,7 +182,8 @@ inline constexpr use_root_sync_key::value_t use_root_sync; === The `root_group` class The `root_group` class implements all member functions common to the -`sycl::group` and `sycl::sub_group` classes. +`sycl::group` and `sycl::sub_group` classes and also contains own +additional functions. [source,c++] ---- @@ -191,6 +192,13 @@ namespace ext { namespace oneapi { namespace experimental { +enum class execution_scope { + work_item, + sub_group, + work_group, + root_group, +}; + template class root_group { public: @@ -221,6 +229,31 @@ public: bool leader() const; + template + std::enable_if_t<(Scope == execution_scope::work_item || + Scope == execution_scope::work_group), + id> + get_id() const; + + template + std::enable_if_t> get_id() const; + + template + size_t get_linear_id() const; + + template + std::enable_if_t<(Scope == execution_scope::work_item || + Scope == execution_scope::work_group), + range> + get_range() const; + + template + std::enable_if_t> + get_range() const; + + template + size_t get_linear_range() const; + }; } // namespace experimental @@ -307,6 +340,67 @@ work-item is the leader of the root-group, and `false` for all other work-items in the root-group. The leader of the root-group is guaranteed to be the work-item for which `get_local_id()` returns 0. +[source,c++] +---- +template +std::enable_if_t<(Scope == execution_scope::work_item || + Scope == execution_scope::work_group), + id> +get_id() const; +---- +_Returns_: An `id` representing the index of the current work-group or work-item at `Scope` +hierarchy level within the `root_group` object. + +[source,c++] +---- +template +std::enable_if_t> get_id() const; +---- +_Returns_: An `id` representing the index of the current sub-group within the +`root_group` object. + +[source,c++] +---- +template +size_t get_linear_id() const; +---- +_Constraints_: `Scope` must be narrower than +`execution_scope::root_group`. + +_Returns_: A linearized number of the current work-group or work-item at `Scope` hierarchy +level within the `root_group` object. + +[source,c++] +---- +template +std::enable_if_t<(Scope == execution_scope::work_item || + Scope == execution_scope::work_group), + range> +get_range() const; +---- +_Returns_: A `range` representing the number of work-groups or work-items of `Scope` +hierarchy level within the `root_group` object. + +[source,c++] +---- +template +std::enable_if_t> +get_range() const; +---- +_Returns_: A `range` representing the number of sub-groups within the `root_group` +object. + +[source,c++] +---- +template +size_t get_linear_range() const; +---- +_Constraints_: `Scope` must be narrower than +`execution_scope::root_group`. + +_Returns_: The number of work-groups or work-items of `Scope` hierarchy level within the +`root_group` object. + === Using a `root_group`