diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index d609027b5816f..20942385c624a 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/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"}; diff --git a/sycl/test-e2e/ESIMD/group_barrier.cpp b/sycl/test-e2e/ESIMD/group_barrier.cpp new file mode 100644 index 0000000000000..492e76225baa1 --- /dev/null +++ b/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 +#include + +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 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 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 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; +} diff --git a/sycl/test/check_device_code/esimd/root_group_barrier.cpp b/sycl/test/check_device_code/esimd/root_group_barrier.cpp new file mode 100644 index 0000000000000..61547a6621054 --- /dev/null +++ b/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 +#include +#include +#include + +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); +}