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

[Wenju He via llvm-branch-commits]

================
@@ -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:

> I don't think this is actually something that should differ between targets. And I'm not sure how to interpret a memory order argument that you need to split. It would add many untested paths

thanks, makes sense to me.

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