diff --git a/adoc/extensions/index.adoc b/adoc/extensions/index.adoc
index a12d8f3fd..71dac4365 100644
--- a/adoc/extensions/index.adoc
+++ b/adoc/extensions/index.adoc
@@ -12,3 +12,4 @@ specification, but their design is subject to change.
 
 include::sycl_khr_default_context.adoc[leveloffset=2]
 include::sycl_khr_queue_empty_query.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 000000000..71b541703
--- /dev/null
+++ b/adoc/extensions/sycl_khr_group_interface.adoc
@@ -0,0 +1,852 @@
+[[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 [code]#sycl::group# and
+[code]#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 a SYCL implementation
+conforms to {cpp23} or later.
+
+[[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 [code]#khr::work_group#, [code]#khr::sub_group# and [code]#khr::member_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
+three types.
+When the synopsis shows an ellipsis ([code]#+/*...*/+#), the subsequent sections
+clarify the definition for each type.
+
+Whether it is possible to query certain properties of a group at compile-time
+depends on the group's type.
+A member function that is not declared [code]#constexpr# in the synopsis below
+may still be declared [code]#constexpr# for a specific group type defined in a
+subsequent section.
+
+[source,role=synopsis]
+----
+
+namespace sycl::khr {
+
+class __group__ {
+
+ public:
+  using id_type = /* ... */;
+  using linear_id_type = /* ... */;
+  using range_type = /* ... */;
+  using extents_type = /* ... */; // C++23
+  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;
+
+  extents_type extents() const noexcept; // C++23
+  extents_type::index_type extent(extents_type::rank_type r) const noexcept; // C++23
+
+  constexpr static extents_type::rank_type rank() noexcept; // C++23
+  constexpr static extents_type::rank_type rank_dynamic() noexcept; // C++23
+  constexpr static size_t static_extent(rank_type r) noexcept; // C++23
+
+  size_type size() const noexcept;
+
+};
+
+template <typename Group>
+member_item<Group> get_member_item(const Group& g) noexcept;
+
+template <typename Group>
+bool leader_of(const Group& g) noexcept;
+
+} // namespace sycl::khr
+----
+
+[[sec:khr-group-interface-common-member-funcs]]
+=== Member functions
+
+.[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
+[api]#+__group__::range+#.
+
+'''
+
+.[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 <<sec:multi-dim-linearization>>) of this
+group within the index space returned by [api]#+__group__::range+#.
+
+'''
+
+.[apidef]#+__group__::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]#+__group__::extents+#
+[source,role=synopsis,id=api:khr-group-interface-common-group-extents]
+----
+constexpr extents_type extents() const noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_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 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.
+
+'''
+
+.[apidef]#+__group__::rank+#
+[source,role=synopsis,id=api:khr-group-interface-common-group-rank]
+----
+static constexpr extents_type::rank_type rank() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_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 extents_type::rank_type rank_dynamic() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_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(extents_type::rank_type r) noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::static_extent(r);#.
+
+'''
+
+.[apidef]#+__group__::size+#
+[source,role=synopsis,id=api:common-group-size]
+----
+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]#khr::get_member_item#
+[source,role=synopsis,id=api:common-group-get-member-item]
+----
+template <typename Group>
+member_item<Group> get_member_item(const Group& g) noexcept;
+----
+
+_Constraints_: [code]#Group# is [code]#work_group# or [code]#sub_group#.
+
+_Returns_: A [code]#member_item# representing the calling work-item within group
+[code]#g#.
+
+'''
+
+.[apidef]#khr::leader_of#
+[source,role=synopsis,id=api:common-group-leader_of]
+----
+template <typename Group>
+bool leader_of(const Group& g) noexcept;
+----
+
+_Constraints_: [code]#Group# is [code]#work_group#, [code]#sub_group# or
+[code]#member_item#.
+
+_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 specific <<work-group>> 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 <<sec:byval-semantics>>) and the common group interface (see
+<<sec:khr-group-interface-common>>).
+
+[source,role=synopsis]
+----
+namespace sycl::khr {
+
+template <int Dimensions = 1>
+class work_group {
+
+ public:
+  using id_type = sycl::id<Dimensions>;
+  using linear_id_type = size_t;
+  using range_type = sycl::range<Dimensions>;
+  using extents_type = std::dextents<size_t, Dimensions>; // C++23
+  using size_type = size_t;
+  static constexpr int dimensions = Dimensions;
+  static constexpr memory_scope fence_scope = memory_scope::work_group;
+
+  work_group(const group<Dimensions>& g) noexcept;
+
+  operator group<Dimensions>() const noexcept;
+
+  /* -- common by-value interface members -- */
+
+  id_type id() const noexcept;
+  linear_id_type linear_id() const noexcept;
+
+  range_type range() const noexcept;
+
+  extents_type extents() const noexcept; // C++23
+  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
+
+  size_type size() const noexcept;
+
+};
+
+} // namespace sycl::khr
+----
+
+.[apititle]#work_group constructor#
+[source,role=synopsis,id=api:khr-group-interface-work-group-constructor]
+----
+work_group(const group<Dimensions>& g) noexcept;
+----
+
+_Effects_: Constructs a [code]#work_group# representing the same collection of
+work-items as [code]#g#.
+
+'''
+
+.[apititle]#work_group conversion operator#
+[source,role=synopsis,id=api:khr-group-interface-work-group-conversion-operator]
+----
+operator group<Dimensions>() const noexcept;
+----
+
+_Returns_: A [code]#group# representing the same collection of work-items as
+this [code]#work_group#.
+
+'''
+
+.[apidef]#+work_group::id+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-id]
+----
+id_type id() const noexcept;
+----
+
+_Returns_: The index of this work-group within the <<nd-range>>.
+
+'''
+
+.[apidef]#+work_group::linear_id+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-linear-id]
+----
+linear_id_type linear_id() const noexcept;
+----
+
+_Returns_: The linearized index (see <<sec:multi-dim-linearization>>) of this
+work-group within the <<nd-range>>.
+
+'''
+
+.[apidef]#+work_group::range+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-range]
+----
+range_type range() const noexcept;
+----
+
+_Returns_: An index space representing all work-groups in the <<nd-range>>.
+
+'''
+
+.[apidef]#+work_group::extents+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-extents]
+----
+extents_type extents() const noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Returns_: The number of work-items in each dimension of the work-group.
+
+'''
+
+.[apidef]#+work_group::extent+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-extent]
+----
+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
+work-group.
+
+'''
+
+.[apidef]#+work_group::rank+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-rank]
+----
+static constexpr extents_type::rank_type rank() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::rank();#.
+
+'''
+
+.[apidef]#+work_group::rank_dynamic+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-rank_dynamic]
+----
+static constexpr extents_type::rank_type rank_dynamic() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#.
+
+'''
+
+.[apidef]#+work_group::static_extent+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-static_extent]
+----
+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);#.
+
+'''
+
+.[apidef]#+work_group::size+#
+[source,role=synopsis,id=api:khr-group-interface-work-group-size]
+----
+size_type size() const noexcept;
+----
+
+_Returns_: The total number of work-items in the work-group, equal to the
+product of the number of work-items in each dimension of the 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 specific <<sub-group>> within a <<work-group>>.
+
+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 <<sec:byval-semantics>>) and the common group interface (see
+<<sec:khr-group-interface-common>>).
+
+[source,role=synopsis]
+----
+namespace sycl::khr {
+
+class sub_group {
+
+ public:
+  using id_type = sycl::id<1>;
+  using linear_id_type = uint32_t;
+  using range_type = sycl::range<1>;
+  using extents_type = std::dextents<uint32_t, 1>; // C++23
+  using size_type = uint32_t;
+  static constexpr int dimensions = 1;
+  static constexpr memory_scope fence_scope = memory_scope::sub_group;
+
+  sub_group(const sycl::sub_group& sg) noexcept;
+
+  operator sycl::sub_group() const noexcept;
+
+  /* -- common by-value interface members -- */
+
+  id_type id() const noexcept;
+  linear_id_type linear_id() const noexcept;
+
+  range_type range() const noexcept;
+
+  extents_type extents() const noexcept; // C++23
+  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
+
+  size_type size() const noexcept;
+  size_type max_size() const noexcept;
+
+};
+
+} // namespace sycl::khr
+----
+
+.[apititle]#sub_group constructor#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-constructor]
+----
+sub_group(const sycl::sub_group& sg) noexcept;
+----
+
+_Effects_: Constructs a [code]#sub_group# representing the same collection of
+work-items as [code]#sg#.
+
+'''
+
+.[apititle]#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]#+sub_group::id+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-id]
+----
+id_type id() const noexcept;
+----
+
+_Returns_: The index of this sub-group within its parent work-group.
+
+'''
+
+.[apidef]#+sub_group::linear_id+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-linear-id]
+----
+linear_id_type linear_id() const noexcept;
+----
+
+_Returns_: The linearized index (see <<sec:multi-dim-linearization>>) of this
+sub-group within its parent work-group.
+
+'''
+
+.[apidef]#+sub_group::range+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-range]
+----
+range_type range() const noexcept;
+----
+
+_Returns_: An index space representing all sub-groups in the same work-group.
+
+'''
+
+.[apidef]#+sub_group::extents+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-extents]
+----
+extents_type extents() const noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Returns_: The number of work-items in each dimension of the sub-group.
+
+'''
+
+.[apidef]#+sub_group::extent+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-extent]
+----
+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 sub-group.
+
+'''
+
+.[apidef]#+sub_group::rank+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-rank]
+----
+static constexpr extents_type::rank_type rank() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::rank();#.
+
+'''
+
+.[apidef]#+sub_group::rank_dynamic+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-rank_dynamic]
+----
+static constexpr extents_type::rank_type rank_dynamic() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#.
+
+'''
+
+.[apidef]#+sub_group::static_extent+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-static_extent]
+----
+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);#.
+
+'''
+
+.[apidef]#+sub_group::size+#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-size]
+----
+size_type size() const noexcept;
+----
+
+_Returns_: The total number of work-items in the sub-group.
+
+'''
+
+.[apidef]#khr::sub_group::max_size#
+[source,role=synopsis,id=api:khr-group-interface-sub-group-max-size]
+----
+size_type max_size() const noexcept;
+----
+
+_Returns_: The maximum number of work-items permitted in any <<sub-group>> 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
+<<device-compiler>>.
+
+'''
+
+[[sec:khr-group-interface-member_item]]
+== [code]#member_item# class
+
+The [code]#member_item# class template encapsulates all functionality required
+to represent a single <<work-item>> within a specific <<group>> 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]#member_item# class
+template.
+Instances of the [code]#member_item# class template can only be acquired from a
+call to [api]#khr::get_member_item#.
+
+The SYCL [code]#member_item# class template provides common by-value semantics
+(see <<sec:byval-semantics>>) and the common group interface (see
+<<sec:khr-group-interface-common>>).
+
+[source,role=synopsis]
+----
+namespace sycl::khr {
+
+template <typename ParentGroup>
+class member_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;
+  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;
+
+  /* -- common by-value interface members -- */
+
+  id_type id() const noexcept;
+  linear_id_type linear_id() const noexcept;
+
+  range_type range() 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
+
+  constexpr size_type size() const noexcept;
+
+};
+
+} // namespace sycl::khr
+----
+
+.[apidef]#+member_item::id+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-id]
+----
+id_type id() const noexcept;
+----
+
+_Returns_: The index of this member-item within its parent group.
+
+'''
+
+.[apidef]#+member_item::linear_id+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-linear-id]
+----
+linear_id_type linear_id() const noexcept;
+----
+
+_Returns_: The linearized index (see <<sec:multi-dim-linearization>>) of this
+member-item within its parent group.
+
+'''
+
+.[apidef]#+member_item::range+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-range]
+----
+range_type range() const noexcept;
+----
+
+_Returns_: An index space representing all member-items in the parent group.
+
+'''
+
+.[apidef]#+member_item::extents+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-extents]
+----
+constexpr extents_type extents() const noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Returns_: An [code]#extents# where all dimensions are 1.
+
+'''
+
+.[apidef]#+member_item::extent+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-extent]
+----
+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_: Equivalent to [code]#return 1;#.
+
+'''
+
+.[apidef]#+member_item::rank+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-rank]
+----
+static constexpr extents_type::rank_type rank() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::rank();#.
+
+'''
+
+.[apidef]#+member_item::rank_dynamic+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-rank_dynamic]
+----
+static constexpr extents_type::rank_type rank_dynamic() noexcept;
+----
+
+_Minimum C++ Version_: {cpp23}
+
+_Effects_: Equivalent to [code]#return extents_type::rank_dynamic();#.
+
+'''
+
+.[apidef]#+member_item::static_extent+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-static_extent]
+----
+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);#.
+
+'''
+
+.[apidef]#+member_item::size+#
+[source,role=synopsis,id=api:khr-group-interface-member-item-size]
+----
+constexpr size_type size() const noexcept;
+----
+
+_Returns_: Equivalent to [code]#return 1;#.
+
+
+[[sec:khr-group-interface-example]]
+== Example
+
+The example below demonstrates the usage of this extension.
+
+[source,,linenums]
+----
+#include <algorithm>
+#include <iostream>
+#include <numeric>
+#include <sycl/sycl.hpp>
+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<int>(N * M, q);
+  int* out = malloc_shared<int>(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::member_item it = get_member_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;
+      free(in, q);
+      free(out, q);
+      exit(-1);
+    }
+  }
+
+  std::cout << "Good computation!" << std::endl;
+  free(in, q);
+  free(out, q);
+  return 0;
+}
+----