[SYCL][ESIMD] Support root group barriers (#15585)

The required driver isn't available in CI yet, but I manually verified
it.

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -53,6 +53,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::multi_ptr<.+>::.+",
     "^sycl::_V1::nd_item<.+>::.+",
     "^sycl::_V1::group<.+>::.+",
+    "^sycl::_V1::group_barrier<.+>",
     "^sycl::_V1::sub_group::.+",
     "^sycl::_V1::range<.+>::.+",
     "^sycl::_V1::kernel_handler::.+",
@@ -64,9 +65,12 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::operator.+<.+>",
     "^sycl::_V1::ext::oneapi::experimental::properties",
     "^sycl::_V1::ext::oneapi::experimental::detail::ExtractProperties",
+    "^sycl::_V1::ext::oneapi::experimental::root_group<.+>::.+",
+    "^sycl::_V1::ext::oneapi::experimental::this_group<.+>",
     "^sycl::_V1::ext::oneapi::sub_group::.+",
     "^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+",
     "^sycl::_V1::ext::oneapi::experimental::this_sub_group",
+    "^sycl::_V1::ext::oneapi::experimental::this_work_item::get_root_group<.+>",
     "^sycl::_V1::ext::oneapi::experimental::uniform<.+>::.+",
     "^sycl::_V1::ext::oneapi::bfloat16::.+",
     "^sycl::_V1::ext::oneapi::experimental::if_architecture_is"};

sycl/test-e2e/ESIMD/group_barrier.cpp

@@ -0,0 +1,68 @@
+//==----- group_barrier.cpp - ESIMD root group barrier test -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===-----------------------------------------------------------===//
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+// REQUIRES-INTEL-DRIVER: lin: 30751
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "esimd_test_utils.hpp"
+#include <sycl/ext/oneapi/experimental/root_group.hpp>
+#include <sycl/group_barrier.hpp>
+
+static constexpr int WorkGroupSize = 16;
+
+static constexpr int VL = 16;
+int main() {
+  bool Pass = true;
+  sycl::queue q;
+  esimd_test::printTestLabel(q);
+  const auto MaxWGs = 8;
+  size_t WorkItemCount = MaxWGs * WorkGroupSize * VL;
+
+  const auto Props = sycl::ext::oneapi::experimental::properties{
+      sycl::ext::oneapi::experimental::use_root_sync};
+  sycl::buffer<int> DataBuf{sycl::range{WorkItemCount}};
+  const auto Range = sycl::nd_range<1>{MaxWGs * WorkGroupSize, WorkGroupSize};
+  q.submit([&](sycl::handler &h) {
+     sycl::accessor Data{DataBuf, h};
+     h.parallel_for(Range, Props, [=](sycl::nd_item<1> it) SYCL_ESIMD_KERNEL {
+       int ID = it.get_global_linear_id();
+       __ESIMD_NS::simd<int, VL> V(ID, 1);
+       // Write data to another kernel's data to verify the barrier works.
+       __ESIMD_NS::block_store(
+           Data, (WorkItemCount * sizeof(int)) - (ID * sizeof(int) * VL), V);
+       if (ID % 2 == 1) {
+         auto Root = it.ext_oneapi_get_root_group();
+         sycl::group_barrier(Root);
+       } else {
+         auto Root =
+             sycl::ext::oneapi::experimental::this_work_item::get_root_group<
+                 1>();
+         sycl::group_barrier(Root);
+       }
+       __ESIMD_NS::simd<int, VL> VOther(ID * VL, 1);
+       __ESIMD_NS::block_store(Data, ID * sizeof(int) * VL, VOther);
+     });
+   }).wait();
+  sycl::host_accessor Data{DataBuf};
+  int ErrCnt = 0;
+  for (int I = 0; I < WorkItemCount; I++) {
+    if (Data[I] != I) {
+      Pass = false;
+      if (++ErrCnt < 16)
+        std::cout << "Data[" << std::to_string(I)
+                  << "] != " << std::to_string(I) << "\n";
+    }
+  }
+  if (Pass)
+    std::cout << "Passed\n";
+  else
+    std::cout << "Failed\n";
+  return !Pass;
+}

sycl/test/check_device_code/esimd/root_group_barrier.cpp

@@ -0,0 +1,15 @@
+// RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o - | FileCheck %s
+
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/ext/oneapi/experimental/root_group.hpp>
+#include <sycl/group_barrier.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::intel::esimd;
+
+SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void
+func(sycl::ext::oneapi::experimental::root_group<1> &rg) {
+  // CHECK: call spir_func void @_Z22__spirv_ControlBarrier{{.*}}(i32 noundef 1, i32 noundef 1, i32 noundef 912)
+  sycl::group_barrier(rg);
+}