[SYCL] Fix sycl::sub_group by-value semantics (#9984)

According to SYCL2020 sycl::sub_group must follow by-value semantics.

Author: KornevNikita

sycl/include/sycl/ext/oneapi/sub_group.hpp

@@ -769,6 +769,25 @@ struct sub_group {
 #endif
   }
 
+  // Common member functions for by-value semantics
+  friend bool operator==(const sub_group &lhs, const sub_group &rhs) {
+#ifdef __SYCL_DEVICE_ONLY__
+    return lhs.get_group_id() == rhs.get_group_id();
+#else
+    throw runtime_error("Sub-groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  friend bool operator!=(const sub_group &lhs, const sub_group &rhs) {
+#ifdef __SYCL_DEVICE_ONLY__
+    return !(lhs == rhs);
+#else
+    throw runtime_error("Sub-groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
 protected:
   template <int dimensions> friend class sycl::nd_item;
   friend sub_group this_sub_group();

sycl/test-e2e/SubGroup/sub_group_by_value_semantics.cpp

@@ -0,0 +1,41 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl.hpp>
+
+int main() {
+  bool result = true;
+  sycl::queue queue;
+  {
+    sycl::buffer<bool, 1> res_buf(&result, 1);
+    queue.submit([&](sycl::handler &cgh) {
+      auto res_acc = res_buf.get_access<sycl::access_mode::read_write>(cgh);
+
+      cgh.parallel_for<class kernel>(sycl::nd_range<3>({1, 1, 1}, {1, 1, 1}),
+                                     [=](sycl::nd_item<3> item) {
+                                       sycl::sub_group a = item.get_sub_group();
+
+                                       // check for reflexivity
+                                       res_acc[0] &= (a == a);
+                                       res_acc[0] &= !(a != a);
+
+                                       // check for symmetry
+                                       auto copied = a;
+                                       auto &b = copied;
+                                       res_acc[0] &= (a == b);
+                                       res_acc[0] &= (b == a);
+                                       res_acc[0] &= !(a != b);
+                                       res_acc[0] &= !(b != a);
+
+                                       // check for transitivity
+                                       auto copiedTwice = copied;
+                                       const auto &c = copiedTwice;
+                                       res_acc[0] &= (c == a);
+                                     });
+    });
+    queue.wait_and_throw();
+  }
+
+  assert(result);
+  return 0;
+}