[llvm-branch-commits] [libclc] libclc: Move sub_group_barrier to clc (PR #185208)
Wenju He via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sat Mar 7 15:58:24 PST 2026
================
@@ -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);
----------------
wenju-he wrote:
our target uses __ATOMIC_SEQ_CST memory order. Could we add `int memory_order` to the parameters of `__clc_sub_group_barrier`?
https://github.com/llvm/llvm-project/pull/185208
More information about the llvm-branch-commits
mailing list