[llvm] [libsycl] Add group_barrier function. (PR #218971)

[Sergey Semenov via llvm-commits]

================
@@ -0,0 +1,78 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <algorithm>
+#include <iostream>
+
+constexpr int LocalSize = 8;
+constexpr int WorkGroups = 4;
+constexpr int GlobalSize = WorkGroups * LocalSize;
+
+static bool runBarrierCase(sycl::queue &Q, int GroupBias, bool MaxCase) {
+  int *Data = sycl::malloc_shared<int>(GlobalSize, Q);
+  int *LocalData = sycl::malloc_shared<int>(GlobalSize, Q);
+
+  Q.parallel_for<class barrier_kernel>(
+      sycl::nd_range<1>{GlobalSize, LocalSize}, [=](sycl::nd_item<1> It) {
+        const int Lid = It.get_local_id(0);
+        const int Gid = It.get_group().get_group_id(0);
+        int *GroupData = LocalData + Gid * LocalSize;
+        GroupData[Lid] = Gid * GroupBias + Lid + 1;
+        sycl::group_barrier(It.get_group());
+
+        if (Lid == 0) {
+          int Result = GroupData[0];
+          if (MaxCase) {
+            for (int I = 1; I < LocalSize; ++I)
+              Result = std::max(Result, GroupData[I]);
+          } else {
+            for (int I = 1; I < LocalSize; ++I)
+              Result += GroupData[I];
+          }
----------------
sergey-semenov wrote:

I'd slightly prefer the latter since, IMO, the current code leaves a false impression that these two scenarios are meaningfully different.

https://github.com/llvm/llvm-project/pull/218971