[llvm-branch-commits] [libclc] libclc: Move sub_group_barrier to clc (PR #185208)

Sat Mar 7 23:53:01 PST 2026

https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/185208

>From a7f32d05755e57b7ffd061af296f7f35b8a3a0f9 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sat, 7 Mar 2026 18:03:47 +0100
Subject: [PATCH] libclc: Move sub_group_barrier to clc

---
 .../synchronization/clc_sub_group_barrier.h   | 21 +++++++++++++++++++
 libclc/clc/lib/amdgcn/SOURCES                 |  1 +
 .../synchronization/clc_sub_group_barrier.cl  | 18 ++++++++++++++++
 libclc/clc/lib/generic/SOURCES                |  1 +
 .../lib/generic/subgroup/sub_group_barrier.cl | 14 +++++++++++++
 libclc/opencl/lib/amdgcn/SOURCES              |  1 -
 libclc/opencl/lib/generic/SOURCES             |  1 +
 .../synchronization/sub_group_barrier.cl      | 14 +++++++------
 8 files changed, 64 insertions(+), 7 deletions(-)
 create mode 100644 libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h
 create mode 100644 libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl
 create mode 100644 libclc/clc/lib/generic/subgroup/sub_group_barrier.cl
 rename libclc/opencl/lib/{amdgcn => generic}/synchronization/sub_group_barrier.cl (56%)

diff --git a/libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h b/libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h
new file mode 100644
index 0000000000000..58ca3d5b5a028
--- /dev/null
+++ b/libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h
@@ -0,0 +1,21 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_SUBGROUP_CLC_SUB_GROUP_BARRIER_H__
+#define __CLC_SUBGROUP_CLC_SUB_GROUP_BARRIER_H__
+
+#include "clc/internal/clc.h"
+#include "clc/mem_fence/clc_mem_semantic.h"
+
+_CLC_DECL _CLC_OVERLOAD void
+__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics,
+                        int memory_scope);
+_CLC_DECL _CLC_OVERLOAD void
+__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics);
+
+#endif // __CLC_SUBGROUP_CLC_SUB_GROUP_BARRIER_H__
diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
index 28ce5f1943825..a280461b1664a 100644
--- a/libclc/clc/lib/amdgcn/SOURCES
+++ b/libclc/clc/lib/amdgcn/SOURCES
@@ -2,6 +2,7 @@ address_space/qualifier.cl
 math/clc_ldexp.cl
 mem_fence/clc_mem_fence.cl
 subgroup/sub_group_broadcast.cl
+synchronization/clc_sub_group_barrier.cl
 synchronization/clc_work_group_barrier.cl
 workitem/clc_get_enqueued_local_size.cl
 workitem/clc_get_global_offset.cl
diff --git a/libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl b/libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl
new file mode 100644
index 0000000000000..ae26b1125f5c0
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl
@@ -0,0 +1,18 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/mem_fence/clc_mem_fence.h"
+#include "clc/synchronization/clc_sub_group_barrier.h"
+
+_CLC_DEF _CLC_OVERLOAD void
+__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics, int scope) {
+  __builtin_amdgcn_wave_barrier();
+
+  if (memory_semantics)
+    __clc_mem_fence(scope, __ATOMIC_ACQ_REL, memory_semantics);
+}
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index 926efe9c6f188..b7ca616299556 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -176,6 +176,7 @@ shared/clc_min.cl
 shared/clc_qualifier.cl
 shared/clc_vload.cl
 shared/clc_vstore.cl
+subgroup/sub_group_barrier.cl
 workitem/clc_get_global_id.cl
 workitem/clc_get_global_linear_id.cl
 workitem/clc_get_local_linear_id.cl
diff --git a/libclc/clc/lib/generic/subgroup/sub_group_barrier.cl b/libclc/clc/lib/generic/subgroup/sub_group_barrier.cl
new file mode 100644
index 0000000000000..45f0a19b9ee39
--- /dev/null
+++ b/libclc/clc/lib/generic/subgroup/sub_group_barrier.cl
@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/synchronization/clc_sub_group_barrier.h"
+
+_CLC_DEF _CLC_OVERLOAD void
+__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics) {
+  __clc_sub_group_barrier(memory_semantics, __MEMORY_SCOPE_WVFRNT);
+}
diff --git a/libclc/opencl/lib/amdgcn/SOURCES b/libclc/opencl/lib/amdgcn/SOURCES
index 8771edfc28f48..7010953d28100 100644
--- a/libclc/opencl/lib/amdgcn/SOURCES
+++ b/libclc/opencl/lib/amdgcn/SOURCES
@@ -1,4 +1,3 @@
 async/wait_group_events.cl
 printf/__printf_alloc.cl
 subgroup/subgroup.cl
-synchronization/sub_group_barrier.cl
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 398e326ffe482..13e8cd97f41c4 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -201,6 +201,7 @@ shared/min.cl
 shared/vload.cl
 shared/vstore.cl
 subgroup/sub_group_broadcast.cl
+synchronization/sub_group_barrier.cl
 synchronization/work_group_barrier.cl
 workitem/get_enqueued_local_size.cl
 workitem/get_global_id.cl
diff --git a/libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl b/libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl
similarity index 56%
rename from libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl
rename to libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl
index 2b57d86294ef5..e03eb01cd767f 100644
--- a/libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl
+++ b/libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl
@@ -6,16 +6,18 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/opencl-base.h>
+#include "clc/opencl/synchronization/utils.h"
+#include "clc/opencl/utils.h"
+#include "clc/synchronization/clc_sub_group_barrier.h"
 
 _CLC_DEF _CLC_OVERLOAD void sub_group_barrier(cl_mem_fence_flags flags,
                                               memory_scope scope) {
-  __builtin_amdgcn_wave_barrier();
-
-  if (flags)
-    atomic_work_item_fence(flags, memory_order_acq_rel, scope);
+  __CLC_MemorySemantics memory_semantics = __opencl_get_memory_semantics(flags);
+  __clc_sub_group_barrier(memory_semantics,
+                          __opencl_get_clang_memory_scope(scope));
 }
 
 _CLC_DEF _CLC_OVERLOAD void sub_group_barrier(cl_mem_fence_flags flags) {
-  sub_group_barrier(flags, memory_scope_sub_group);
+  __CLC_MemorySemantics memory_semantics = __opencl_get_memory_semantics(flags);
+  __clc_sub_group_barrier(memory_semantics);
 }

