[llvm] [libsycl] add group_barrier function (PR #218971)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 9 08:47:39 PDT 2026
https://github.com/KseniyaTikhomirova updated https://github.com/llvm/llvm-project/pull/218971
>From 82ac329886575702d8ff137fad656ecb99c5506c Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 26 Aug 2026 05:03:27 -0700
Subject: [PATCH 1/2] add group_barrier function
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/include/sycl/__impl/group_barrier.hpp | 85 +++++++++++++++++++
libsycl/include/sycl/__spirv/spirv_types.hpp | 45 ++++++++++
libsycl/include/sycl/sycl.hpp | 1 +
libsycl/test/basic/group_barrier.cpp | 78 +++++++++++++++++
4 files changed, 209 insertions(+)
create mode 100644 libsycl/include/sycl/__impl/group_barrier.hpp
create mode 100644 libsycl/include/sycl/__spirv/spirv_types.hpp
create mode 100644 libsycl/test/basic/group_barrier.cpp
diff --git a/libsycl/include/sycl/__impl/group_barrier.hpp b/libsycl/include/sycl/__impl/group_barrier.hpp
new file mode 100644
index 0000000000000..b03206c60b96e
--- /dev/null
+++ b/libsycl/include/sycl/__impl/group_barrier.hpp
@@ -0,0 +1,85 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the SYCL 2020 group_barrier function (4.17.2.3.).
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_GROUP_BARRIER_HPP
+#define _LIBSYCL___IMPL_GROUP_BARRIER_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/group.hpp>
+#include <sycl/__impl/memory_enums.hpp>
+#include <sycl/__impl/sub_group.hpp>
+#include <sycl/__spirv/spirv_types.hpp>
+
+#include <type_traits>
+
+void __spirv_ControlBarrier(std::uint32_t Execution, std::uint32_t Memory,
+ std::uint32_t Semantics);
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+// 4.17.2.1. Group type trait.
+
+template <class T> struct is_group : std::false_type {};
+
+template <int Dimensions>
+struct is_group<group<Dimensions>> : std::true_type {};
+
+template <> struct is_group<sub_group> : std::true_type {};
+
+template <class T>
+inline constexpr bool is_group_v = is_group<std::decay_t<T>>::value;
+
+namespace detail {
+
+static constexpr __spirv::Scope getScope(memory_scope Scope) {
+ switch (Scope) {
+ case memory_scope::work_item:
+ return __spirv::Scope::Invocation;
+ case memory_scope::sub_group:
+ return __spirv::Scope::Subgroup;
+ case memory_scope::work_group:
+ return __spirv::Scope::Workgroup;
+ case memory_scope::device:
+ return __spirv::Scope::Device;
+ case memory_scope::system:
+ return __spirv::Scope::CrossDevice;
+ }
+}
+
+template <typename Group> struct group_scope {};
+
+template <int Dimensions> struct group_scope<group<Dimensions>> {
+ static constexpr __spirv::Scope value = __spirv::Scope::Workgroup;
+};
+
+template <> struct group_scope<::sycl::sub_group> {
+ static constexpr __spirv::Scope value = __spirv::Scope::Subgroup;
+};
+
+} // namespace detail
+
+/// Blocks until all work-items in group g have reached this synchronization
+/// point.
+template <typename Group>
+std::enable_if_t<is_group_v<Group>>
+group_barrier(Group /*G*/, memory_scope FenceScope = Group::fence_scope) {
+ __spirv_ControlBarrier(detail::group_scope<Group>::value,
+ detail::getScope(FenceScope),
+ __spirv::MemorySemanticsMask::SequentiallyConsistent |
+ __spirv::MemorySemanticsMask::SubgroupMemory |
+ __spirv::MemorySemanticsMask::WorkgroupMemory);
+}
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_GROUP_BARRIER_HPP
diff --git a/libsycl/include/sycl/__spirv/spirv_types.hpp b/libsycl/include/sycl/__spirv/spirv_types.hpp
new file mode 100644
index 0000000000000..47f29367df3b4
--- /dev/null
+++ b/libsycl/include/sycl/__spirv/spirv_types.hpp
@@ -0,0 +1,45 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains common SPIR-V enum definitions used by libsycl.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___SPIRV_SPIRV_TYPES_HPP
+#define _LIBSYCL___SPIRV_SPIRV_TYPES_HPP
+
+#include <cstdint>
+
+namespace __spirv {
+
+enum Scope : int32_t {
+ CrossDevice = 0,
+ Device = 1,
+ Workgroup = 2,
+ Subgroup = 3,
+ Invocation = 4,
+};
+
+enum MemorySemanticsMask : int32_t {
+ None = 0x0,
+ Acquire = 0x2,
+ Release = 0x4,
+ AcquireRelease = 0x8,
+ SequentiallyConsistent = 0x10,
+ UniformMemory = 0x40,
+ SubgroupMemory = 0x80,
+ WorkgroupMemory = 0x100,
+ CrossWorkgroupMemory = 0x200,
+ AtomicCounterMemory = 0x400,
+ ImageMemory = 0x800,
+};
+
+} // namespace __spirv
+
+#endif // _LIBSYCL___SPIRV_SPIRV_TYPES_HPP
diff --git a/libsycl/include/sycl/sycl.hpp b/libsycl/include/sycl/sycl.hpp
index 597c3377be27d..4c5e3e62a11c8 100644
--- a/libsycl/include/sycl/sycl.hpp
+++ b/libsycl/include/sycl/sycl.hpp
@@ -20,6 +20,7 @@
#include <sycl/__impl/event.hpp>
#include <sycl/__impl/exception.hpp>
#include <sycl/__impl/group.hpp>
+#include <sycl/__impl/group_barrier.hpp>
#include <sycl/__impl/index_space_classes.hpp>
#include <sycl/__impl/memory_enums.hpp>
#include <sycl/__impl/nd_item.hpp>
diff --git a/libsycl/test/basic/group_barrier.cpp b/libsycl/test/basic/group_barrier.cpp
new file mode 100644
index 0000000000000..e4c964f70bbde
--- /dev/null
+++ b/libsycl/test/basic/group_barrier.cpp
@@ -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];
+ }
+ GroupData[0] = Result;
+ }
+
+ sycl::group_barrier(It.get_group());
+ Data[It.get_global_id(0)] = GroupData[0];
+ });
+
+ Q.wait();
+
+ bool Failure = false;
+ for (int Gid = 0; Gid < WorkGroups; ++Gid) {
+ int Expected = 0;
+ for (int Lid = 0; Lid < LocalSize; ++Lid) {
+ const int Value = Gid * GroupBias + Lid + 1;
+ Expected += Value;
+ }
+ if (MaxCase)
+ Expected = Gid * GroupBias + LocalSize;
+
+ for (int Lid = 0; Lid < LocalSize; ++Lid) {
+ const int Index = Gid * LocalSize + Lid;
+ if (Data[Index] != Expected) {
+ std::cerr << "Mismatch at group " << Gid << " lane " << Lid << ": got "
+ << Data[Index] << ", expected " << Expected << std::endl;
+ Failure = true;
+ }
+ }
+ }
+
+ sycl::free(Data, Q);
+ sycl::free(LocalData, Q);
+ return !Failure;
+}
+
+int main() {
+ sycl::queue Q;
+
+ if (!runBarrierCase(Q, 10, false))
+ return 1;
+ if (!runBarrierCase(Q, 17, true))
+ return 1;
+
+ return 0;
+}
>From 0a26e7f289c8a879d450668169d0e7b21be52c3a Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Wed, 9 Sep 2026 08:47:23 -0700
Subject: [PATCH 2/2] fix code-review comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/include/sycl/__impl/group_barrier.hpp | 1 +
.../test/basic/group_barrier_device_code.cpp | 38 +++++++++++++++++++
2 files changed, 39 insertions(+)
create mode 100644 libsycl/test/basic/group_barrier_device_code.cpp
diff --git a/libsycl/include/sycl/__impl/group_barrier.hpp b/libsycl/include/sycl/__impl/group_barrier.hpp
index b03206c60b96e..d23fb0294616f 100644
--- a/libsycl/include/sycl/__impl/group_barrier.hpp
+++ b/libsycl/include/sycl/__impl/group_barrier.hpp
@@ -20,6 +20,7 @@
#include <sycl/__impl/sub_group.hpp>
#include <sycl/__spirv/spirv_types.hpp>
+#include <cstdint>
#include <type_traits>
void __spirv_ControlBarrier(std::uint32_t Execution, std::uint32_t Memory,
diff --git a/libsycl/test/basic/group_barrier_device_code.cpp b/libsycl/test/basic/group_barrier_device_code.cpp
new file mode 100644
index 0000000000000..4387c77c01373
--- /dev/null
+++ b/libsycl/test/basic/group_barrier_device_code.cpp
@@ -0,0 +1,38 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
+
+// Checks the execution scope, memory scope and memory semantics operands that
+// group_barrier passes to __spirv_ControlBarrier.
+
+#include <sycl/sycl.hpp>
+
+void test(sycl::queue Q) {
+ Q.parallel_for<class barrier_kernel>(
+ sycl::nd_range<1>{8, 4}, [=](sycl::nd_item<1> It) {
+ // Execution scope is Workgroup (2), memory scope defaults to the
+ // group's fence_scope, i.e. Workgroup (2).
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 2, i32 noundef 2, i32 noundef 400)
+ sycl::group_barrier(It.get_group());
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 2, i32 noundef 4, i32 noundef 400)
+ sycl::group_barrier(It.get_group(), sycl::memory_scope::work_item);
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 2, i32 noundef 3, i32 noundef 400)
+ sycl::group_barrier(It.get_group(), sycl::memory_scope::sub_group);
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 2, i32 noundef 2, i32 noundef 400)
+ sycl::group_barrier(It.get_group(), sycl::memory_scope::work_group);
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 2, i32 noundef 1, i32 noundef 400)
+ sycl::group_barrier(It.get_group(), sycl::memory_scope::device);
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 2, i32 noundef 0, i32 noundef 400)
+ sycl::group_barrier(It.get_group(), sycl::memory_scope::system);
+
+ // Execution scope is Subgroup (3), memory scope defaults to the
+ // sub-group's fence_scope, i.e. Subgroup (3).
+ // CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef
+ // 3, i32 noundef 3, i32 noundef 400)
+ sycl::group_barrier(It.get_sub_group());
+ });
+}
More information about the llvm-commits
mailing list