[libclc] 4c9d448 - libclc: Move sub_group_barrier to clc (#185208)
via cfe-commits
cfe-commits at lists.llvm.org
Sun Mar 8 00:09:44 PST 2026
Author: Matt Arsenault
Date: 2026-03-08T09:09:40+01:00
New Revision: 4c9d448e058d2a6371634a27608ef3fff01e817d
URL: https://github.com/llvm/llvm-project/commit/4c9d448e058d2a6371634a27608ef3fff01e817d
DIFF: https://github.com/llvm/llvm-project/commit/4c9d448e058d2a6371634a27608ef3fff01e817d.diff
LOG: libclc: Move sub_group_barrier to clc (#185208)
Added:
libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h
libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl
libclc/clc/lib/generic/subgroup/sub_group_barrier.cl
libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl
Modified:
libclc/clc/lib/amdgcn/SOURCES
libclc/clc/lib/generic/SOURCES
libclc/opencl/lib/amdgcn/SOURCES
libclc/opencl/lib/generic/SOURCES
Removed:
libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl
################################################################################
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 c29f79b82e882..8e2df4a3e910a 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -202,6 +202,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);
}
More information about the cfe-commits
mailing list