[SYCL] Add support for SYCL 2020 in class group

sycl/include/CL/sycl/group.hpp

@@ -104,53 +104,67 @@ template <int Dimensions = 1> class group {
 
   group() = delete;
 
+  __SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead")
   id<Dimensions> get_id() const { return index; }
 
+  __SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead")
   size_t get_id(int dimension) const { return index[dimension]; }
 
+  id<Dimensions> get_group_id() const { return index; }
+
+  size_t get_group_id(int dimension) const { return index[dimension]; }
+
   range<Dimensions> get_global_range() const { return globalRange; }
 
   size_t get_global_range(int dimension) const {
     return globalRange[dimension];
   }
 
+  id<Dimensions> get_local_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>();
+#else
+    throw runtime_error("get_local_id() is not implemented on host device",
+                        PI_INVALID_DEVICE);
+    // Implementing get_local_id() on host device requires ABI breaking change.
+    // It requires extending class group with local item which represents
+    // local_id. Currently this local id is only used in nd_item and group
+    // cannot access it.
+#endif
+  }
+
+  size_t get_local_linear_id() const {
+    return get_local_linear_id_impl<Dimensions>();
+  }
+
   range<Dimensions> get_local_range() const { return localRange; }
 
   size_t get_local_range(int dimension) const { return localRange[dimension]; }
 
+  size_t get_local_linear_range() const {
+    return get_local_linear_range_impl();
+  }
+
   range<Dimensions> get_group_range() const { return groupRange; }
 
   size_t get_group_range(int dimension) const {
     return get_group_range()[dimension];
   }
 
+  size_t get_group_linear_range() const {
+    return get_group_linear_range_impl();
+  }
+
+  range<Dimensions> get_max_local_range() const { return get_local_range(); }
+
   size_t operator[](int dimension) const { return index[dimension]; }
 
-  template <int dims = Dimensions>
-  typename detail::enable_if_t<(dims == 1), size_t> get_linear_id() const {
-    return index[0];
-  }
+  __SYCL2020_DEPRECATED("use sycl::group::get_group_linear_id() instead")
+  size_t get_linear_id() const { return get_group_linear_id(); }
 
-  template <int dims = Dimensions>
-  typename detail::enable_if_t<(dims == 2), size_t> get_linear_id() const {
-    return index[0] * groupRange[1] + index[1];
-  }
+  size_t get_group_linear_id() const { return get_group_linear_id_impl(); }
 
-  // SYCL specification 1.2.1rev5, section 4.7.6.5 "Buffer accessor":
-  //    Whenever a multi-dimensional index is passed to a SYCL accessor the
-  //    linear index is calculated based on the index {id1, id2, id3} provided
-  //    and the range of the SYCL accessor {r1, r2, r3} according to row-major
-  //    ordering as follows:
-  //      id3 + (id2 · r3) + (id1 · r3 · r2)            (4.3)
-  // section 4.8.1.8 "group class":
-  //    size_t get_linear_id()const
-  //    Get a linearized version of the work-group id. Calculating a linear
-  //    work-group id from a multi-dimensional index follows the equation 4.3.
-  template <int dims = Dimensions>
-  typename detail::enable_if_t<(dims == 3), size_t> get_linear_id() const {
-    return (index[0] * groupRange[1] * groupRange[2]) +
-           (index[1] * groupRange[2]) + index[2];
-  }
+  bool leader() const { return (get_local_linear_id() == 0); }
 
   template <typename WorkItemFunctionT>
   void parallel_for_work_item(WorkItemFunctionT Func) const {
@@ -397,6 +411,99 @@ template <int Dimensions = 1> class group {
   range<Dimensions> groupRange;
   id<Dimensions> index;
 
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_local_linear_id_impl() const {
+    id<Dimensions> localId = get_local_id();
+    return localId[0];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 2), size_t>
+  get_local_linear_id_impl() const {
+    id<Dimensions> localId = get_local_id();
+    return localId[0] * groupRange[1] + localId[1];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 3), size_t>
+  get_local_linear_id_impl() const {
+    id<Dimensions> localId = get_local_id();
+    return (localId[0] * groupRange[1] * groupRange[2]) +
+           (localId[1] * groupRange[2]) + localId[2];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_local_linear_range_impl() const {
+    auto localRange = get_local_range();
+    return localRange[0];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 2), size_t>
+  get_local_linear_range_impl() const {
+    auto localRange = get_local_range();
+    return localRange[0] * localRange[1];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 3), size_t>
+  get_local_linear_range_impl() const {
+    auto localRange = get_local_range();
+    return localRange[0] * localRange[1] * localRange[2];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_group_linear_range_impl() const {
+    auto groupRange = get_group_range();
+    return groupRange[0];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 2), size_t>
+  get_group_linear_range_impl() const {
+    auto groupRange = get_group_range();
+    return groupRange[0] * groupRange[1];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 3), size_t>
+  get_group_linear_range_impl() const {
+    auto groupRange = get_group_range();
+    return groupRange[0] * groupRange[1] * groupRange[2];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_group_linear_id_impl() const {
+    return index[0];
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 2), size_t>
+  get_group_linear_id_impl() const {
+    return index[0] * groupRange[1] + index[1];
+  }
+
+  // SYCL specification 1.2.1rev5, section 4.7.6.5 "Buffer accessor":
+  //    Whenever a multi-dimensional index is passed to a SYCL accessor the
+  //    linear index is calculated based on the index {id1, id2, id3} provided
+  //    and the range of the SYCL accessor {r1, r2, r3} according to row-major
+  //    ordering as follows:
+  //      id3 + (id2 · r3) + (id1 · r3 · r2)            (4.3)
+  // section 4.8.1.8 "group class":
+  //    size_t get_linear_id()const
+  //    Get a linearized version of the work-group id. Calculating a linear
+  //    work-group id from a multi-dimensional index follows the equation 4.3.
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 3), size_t>
+  get_group_linear_id_impl() const {
+    return (index[0] * groupRange[1] * groupRange[2]) +
+           (index[1] * groupRange[2]) + index[2];
+  }
+
   void waitForHelper() const {}
 
   void waitForHelper(device_event Event) const {

sycl/test/basic_tests/group.cpp

@@ -29,6 +29,19 @@ int main() {
   assert(one_dim.get_group_range(0) == 2);
   assert(one_dim[0] == 1);
   assert(one_dim.get_linear_id() == 1);
+  assert(one_dim.get_group_linear_id() == 1);
+
+  try {
+    one_dim.get_local_id();
+    assert(0); // get_local_id() is not implemented on host device
+  } catch (cl::sycl::runtime_error) {
+  }
+
+  try {
+    one_dim.get_local_linear_id();
+    assert(0); // get_local_id() is not implemented on host device
+  } catch (cl::sycl::runtime_error) {
+  }
 
   // two dimension group
   cl::sycl::group<2> two_dim = Builder::createGroup<2>({8, 4}, {4, 2}, {1, 1});
@@ -47,6 +60,19 @@ int main() {
   assert(two_dim[0] == 1);
   assert(two_dim[1] == 1);
   assert(two_dim.get_linear_id() == 3);
+  assert(two_dim.get_group_linear_id() == 3);
+
+  try {
+    two_dim.get_local_id();
+    assert(0); // get_local_id() is not implemented on host device
+  } catch (cl::sycl::runtime_error) {
+  }
+
+  try {
+    two_dim.get_local_linear_id();
+    assert(0); // get_local_id() is not implemented on host device
+  } catch (cl::sycl::runtime_error) {
+  }
 
   // three dimension group
   cl::sycl::group<3> three_dim =
@@ -71,4 +97,17 @@ int main() {
   assert(three_dim[1] == 1);
   assert(three_dim[2] == 1);
   assert(three_dim.get_linear_id() == 7);
+  assert(three_dim.get_group_linear_id() == 7);
+
+  try {
+    three_dim.get_local_id();
+    assert(0); // get_local_id() is not implemented on host device
+  } catch (cl::sycl::runtime_error) {
+  }
+
+  try {
+    three_dim.get_local_linear_id();
+    assert(0); // get_local_id() is not implemented on host device
+  } catch (cl::sycl::runtime_error) {
+  }
 }

sycl/test/warnings/sycl_2020_deprecations.cpp

@@ -179,5 +179,14 @@ int main() {
   // expected-warning@+1 {{'atomic<int, sycl::access::address_space::global_space>' is deprecated: sycl::atomic is deprecated since SYCL 2020}}
   cl::sycl::atomic<int> b(a);
 
+  cl::sycl::group<1> group =
+      cl::sycl::detail::Builder::createGroup<1>({8}, {4}, {1});
+  // expected-warning@+1{{'get_id' is deprecated: use sycl::group::get_group_id() instead}}
+  group.get_id();
+  // expected-warning@+1{{'get_id' is deprecated: use sycl::group::get_group_id() instead}}
+  group.get_id(1);
+  // expected-warning@+1{{'get_linear_id' is deprecated: use sycl::group::get_group_linear_id() instead}}
+  group.get_linear_id();
+
   return 0;
 }